From 06547efd90bdb16433f65ccdf31bdef755de5e2a Mon Sep 17 00:00:00 2001 From: Qianfeng Zhang Date: Fri, 12 Jun 2026 05:31:09 +0000 Subject: [PATCH] Remove the kHasBias==true instances to save building time --- .../18_hstu_attention/generate_instances.py | 4 +- .../hstu_attention_group_forward_bf16.cpp | 4 +- .../hstu_attention_group_forward_fp16.cpp | 4 +- .../hstu_attention_no_group_forward_bf16.cpp | 4 +- .../hstu_attention_no_group_forward_fp16.cpp | 4 +- ...lse_false_has_bias_no_dropout_maxk_128.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_256.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_64.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_96.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_128.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_256.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_64.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_96.cpp | 20 -- ..._lse_true_has_bias_no_dropout_maxk_128.cpp | 20 -- ..._lse_true_has_bias_no_dropout_maxk_256.cpp | 20 -- ...e_lse_true_has_bias_no_dropout_maxk_64.cpp | 20 -- ...e_lse_true_has_bias_no_dropout_maxk_96.cpp | 20 -- ...ion_batched_forward_bf16_instances_ref.hpp | 216 ------------------ ...lse_false_has_bias_no_dropout_maxk_128.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_256.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_64.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_96.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_128.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_256.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_64.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_96.cpp | 20 -- ..._lse_true_has_bias_no_dropout_maxk_128.cpp | 20 -- ..._lse_true_has_bias_no_dropout_maxk_256.cpp | 20 -- ...e_lse_true_has_bias_no_dropout_maxk_64.cpp | 20 -- ...e_lse_true_has_bias_no_dropout_maxk_96.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_128.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_256.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_64.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_96.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_128.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_256.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_64.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_96.cpp | 20 -- ..._lse_true_has_bias_no_dropout_maxk_128.cpp | 20 -- ..._lse_true_has_bias_no_dropout_maxk_256.cpp | 20 -- ...e_lse_true_has_bias_no_dropout_maxk_64.cpp | 20 -- ...e_lse_true_has_bias_no_dropout_maxk_96.cpp | 20 -- ...ion_batched_forward_fp16_instances_ref.hpp | 216 ------------------ ...lse_false_has_bias_no_dropout_maxk_128.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_256.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_64.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_96.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_128.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_256.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_64.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_96.cpp | 20 -- ..._lse_true_has_bias_no_dropout_maxk_128.cpp | 20 -- ..._lse_true_has_bias_no_dropout_maxk_256.cpp | 20 -- ...e_lse_true_has_bias_no_dropout_maxk_64.cpp | 20 -- ...e_lse_true_has_bias_no_dropout_maxk_96.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_128.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_256.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_64.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_96.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_128.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_256.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_64.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_96.cpp | 20 -- ..._lse_true_has_bias_no_dropout_maxk_128.cpp | 20 -- ..._lse_true_has_bias_no_dropout_maxk_256.cpp | 20 -- ...e_lse_true_has_bias_no_dropout_maxk_64.cpp | 20 -- ...e_lse_true_has_bias_no_dropout_maxk_96.cpp | 20 -- ...ntion_group_forward_bf16_instances_ref.hpp | 216 ------------------ ...lse_false_has_bias_no_dropout_maxk_128.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_256.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_64.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_96.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_128.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_256.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_64.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_96.cpp | 20 -- ..._lse_true_has_bias_no_dropout_maxk_128.cpp | 20 -- ..._lse_true_has_bias_no_dropout_maxk_256.cpp | 20 -- ...e_lse_true_has_bias_no_dropout_maxk_64.cpp | 20 -- ...e_lse_true_has_bias_no_dropout_maxk_96.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_128.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_256.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_64.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_96.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_128.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_256.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_64.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_96.cpp | 20 -- ..._lse_true_has_bias_no_dropout_maxk_128.cpp | 20 -- ..._lse_true_has_bias_no_dropout_maxk_256.cpp | 20 -- ...e_lse_true_has_bias_no_dropout_maxk_64.cpp | 20 -- ...e_lse_true_has_bias_no_dropout_maxk_96.cpp | 20 -- ...ntion_group_forward_fp16_instances_ref.hpp | 216 ------------------ ...lse_false_has_bias_no_dropout_maxk_128.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_256.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_64.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_96.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_128.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_256.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_64.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_96.cpp | 20 -- ..._lse_true_has_bias_no_dropout_maxk_128.cpp | 20 -- ..._lse_true_has_bias_no_dropout_maxk_256.cpp | 20 -- ...e_lse_true_has_bias_no_dropout_maxk_64.cpp | 20 -- ...e_lse_true_has_bias_no_dropout_maxk_96.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_128.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_256.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_64.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_96.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_128.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_256.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_64.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_96.cpp | 20 -- ..._lse_true_has_bias_no_dropout_maxk_128.cpp | 20 -- ..._lse_true_has_bias_no_dropout_maxk_256.cpp | 20 -- ...e_lse_true_has_bias_no_dropout_maxk_64.cpp | 20 -- ...e_lse_true_has_bias_no_dropout_maxk_96.cpp | 20 -- ...tion_jagged_forward_bf16_instances_ref.hpp | 216 ------------------ ...lse_false_has_bias_no_dropout_maxk_128.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_256.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_64.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_96.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_128.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_256.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_64.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_96.cpp | 20 -- ..._lse_true_has_bias_no_dropout_maxk_128.cpp | 20 -- ..._lse_true_has_bias_no_dropout_maxk_256.cpp | 20 -- ...e_lse_true_has_bias_no_dropout_maxk_64.cpp | 20 -- ...e_lse_true_has_bias_no_dropout_maxk_96.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_128.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_256.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_64.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_96.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_128.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_256.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_64.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_96.cpp | 20 -- ..._lse_true_has_bias_no_dropout_maxk_128.cpp | 20 -- ..._lse_true_has_bias_no_dropout_maxk_256.cpp | 20 -- ...e_lse_true_has_bias_no_dropout_maxk_64.cpp | 20 -- ...e_lse_true_has_bias_no_dropout_maxk_96.cpp | 20 -- ...tion_jagged_forward_fp16_instances_ref.hpp | 216 ------------------ ...lse_false_has_bias_no_dropout_maxk_128.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_256.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_64.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_96.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_128.cpp | 20 -- ...lse_false_has_bias_no_dropout_maxk_256.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_64.cpp | 20 -- ..._lse_false_has_bias_no_dropout_maxk_96.cpp | 20 -- ..._lse_true_has_bias_no_dropout_maxk_128.cpp | 20 -- ..._lse_true_has_bias_no_dropout_maxk_256.cpp | 20 -- ...e_lse_true_has_bias_no_dropout_maxk_64.cpp | 20 -- ...e_lse_true_has_bias_no_dropout_maxk_96.cpp | 20 -- 155 files changed, 10 insertions(+), 4186 deletions(-) delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp diff --git a/example/ck_tile/18_hstu_attention/generate_instances.py b/example/ck_tile/18_hstu_attention/generate_instances.py index f614460ad4..e9b2de3ab2 100644 --- a/example/ck_tile/18_hstu_attention/generate_instances.py +++ b/example/ck_tile/18_hstu_attention/generate_instances.py @@ -99,7 +99,7 @@ def create_forward_instances(instance_dir: Path, headdims: List) -> None: (True, True), (False, False), ]: - for has_bias in [True, False]: + for has_bias in [False]: for has_dropout in [False]: for max_k in headdims: fname = HSTU_FORWARD_INSTANCE_FNAME.format( @@ -158,7 +158,7 @@ def create_forward_instances_ref(instance_dir: Path, headdims: List) -> None: file.write(HSTU_COPYRIGHT_HEADER) file.write(forward_instance_inc) for max_k in headdims: - for has_bias in [True, False]: + for has_bias in [False]: for has_dropout in [False]: for has_causal in [True, False]: for use_softmax, store_lse in [ diff --git a/example/ck_tile/18_hstu_attention/hstu_attention_group_forward_bf16.cpp b/example/ck_tile/18_hstu_attention/hstu_attention_group_forward_bf16.cpp index 1d411126dd..0631b26cbd 100644 --- a/example/ck_tile/18_hstu_attention/hstu_attention_group_forward_bf16.cpp +++ b/example/ck_tile/18_hstu_attention/hstu_attention_group_forward_bf16.cpp @@ -8,11 +8,11 @@ void hstu_attention_group_forward_bf16(HstuAttentionGroupFwdParams& param, hipStream_t stream) { - const bool has_bias = (param.bias_ptr != nullptr); const bool use_causal = param.use_causal; bool store_lse = (param.use_softmax && param.is_training); - BOOL_SWITCH_3(has_bias, kHasBias, use_causal, kUseCausal, param.use_softmax, kUseSoftmax, [&] { + constexpr bool kHasBias = false; + BOOL_SWITCH_2(use_causal, kUseCausal, param.use_softmax, kUseSoftmax, [&] { HDIM_SWITCH(param.hdim_qk, param.hdim_v, MaxK, [&] { BOOL_SWITCH(store_lse, kStoreLSE, [&] { if constexpr(kUseSoftmax || !kStoreLSE) diff --git a/example/ck_tile/18_hstu_attention/hstu_attention_group_forward_fp16.cpp b/example/ck_tile/18_hstu_attention/hstu_attention_group_forward_fp16.cpp index a418a83e00..75d7283560 100644 --- a/example/ck_tile/18_hstu_attention/hstu_attention_group_forward_fp16.cpp +++ b/example/ck_tile/18_hstu_attention/hstu_attention_group_forward_fp16.cpp @@ -8,11 +8,11 @@ void hstu_attention_group_forward_fp16(HstuAttentionGroupFwdParams& param, hipStream_t stream) { - const bool has_bias = (param.bias_ptr != nullptr); const bool use_causal = param.use_causal; const bool store_lse = (param.use_softmax && param.is_training); - BOOL_SWITCH_3(has_bias, kHasBias, use_causal, kUseCausal, param.use_softmax, kUseSoftmax, [&] { + constexpr bool kHasBias = false; + BOOL_SWITCH_2(use_causal, kUseCausal, param.use_softmax, kUseSoftmax, [&] { HDIM_SWITCH(param.hdim_qk, param.hdim_v, MaxK, [&] { BOOL_SWITCH(store_lse, kStoreLSE, [&] { if constexpr(kUseSoftmax || !kStoreLSE) diff --git a/example/ck_tile/18_hstu_attention/hstu_attention_no_group_forward_bf16.cpp b/example/ck_tile/18_hstu_attention/hstu_attention_no_group_forward_bf16.cpp index a28fc15edc..56723230c9 100644 --- a/example/ck_tile/18_hstu_attention/hstu_attention_no_group_forward_bf16.cpp +++ b/example/ck_tile/18_hstu_attention/hstu_attention_no_group_forward_bf16.cpp @@ -9,11 +9,11 @@ void hstu_attention_no_group_forward_bf16(HstuAttentionNoGroupFwdParams& param, hipStream_t stream) { - const bool has_bias = (param.bias_ptr != nullptr); const bool use_causal = param.use_causal; bool store_lse = (param.use_softmax && param.is_training); - BOOL_SWITCH_3(has_bias, kHasBias, use_causal, kUseCausal, param.use_softmax, kUseSoftmax, [&] { + constexpr bool kHasBias = false; + BOOL_SWITCH_2(use_causal, kUseCausal, param.use_softmax, kUseSoftmax, [&] { HDIM_SWITCH(param.hdim_qk, param.hdim_v, MaxK, [&] { BOOL_SWITCH(store_lse, kStoreLSE, [&] { if constexpr(kUseSoftmax || !kStoreLSE) diff --git a/example/ck_tile/18_hstu_attention/hstu_attention_no_group_forward_fp16.cpp b/example/ck_tile/18_hstu_attention/hstu_attention_no_group_forward_fp16.cpp index 06aef656f7..24093880aa 100644 --- a/example/ck_tile/18_hstu_attention/hstu_attention_no_group_forward_fp16.cpp +++ b/example/ck_tile/18_hstu_attention/hstu_attention_no_group_forward_fp16.cpp @@ -9,11 +9,11 @@ void hstu_attention_no_group_forward_fp16(HstuAttentionNoGroupFwdParams& param, hipStream_t stream) { - const bool has_bias = (param.bias_ptr != nullptr); const bool use_causal = param.use_causal; bool store_lse = (param.use_softmax && param.is_training); - BOOL_SWITCH_3(has_bias, kHasBias, use_causal, kUseCausal, param.use_softmax, kUseSoftmax, [&] { + constexpr bool kHasBias = false; + BOOL_SWITCH_2(use_causal, kUseCausal, param.use_softmax, kUseSoftmax, [&] { HDIM_SWITCH(param.hdim_qk, param.hdim_v, MaxK, [&] { BOOL_SWITCH(store_lse, kStoreLSE, [&] { if constexpr(kUseSoftmax || !kStoreLSE) diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_lse_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_lse_false_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index e6d8258e13..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::bf16_t, - true, - false, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_lse_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_lse_false_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index e7e82a177d..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::bf16_t, - true, - false, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_lse_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_lse_false_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 9841c960a1..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::bf16_t, - true, - false, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index 6c5e7f5a11..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::bf16_t, - true, - false, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index 0dccdd8edd..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 38c8ccb2df..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index a998ee426e..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index 22207f382f..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_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_lse_true_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index 88b93bcef4..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_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_lse_true_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index ceaf73d47b..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_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_lse_true_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 99e916571c..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index e6acd7d80f..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git 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 index bdabd98f53..fbacacbfeb 100644 --- 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 @@ -10,60 +10,6 @@ #include "hstu_attention_batched_forward_dispatch.hpp" #include "hstu_attention_params.hpp" -extern template void run_batched_forward_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::bf16_t, - true, - false, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::bf16_t, - false, - true, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::bf16_t, - false, - true, - true, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - extern template void run_batched_forward_dispatch< ck_tile::bf16_t, true, @@ -118,60 +64,6 @@ extern template void run_batched_forward_dispatch< false, 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); -extern template void run_batched_forward_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::bf16_t, - true, - false, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::bf16_t, - false, - true, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::bf16_t, - false, - true, - true, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - extern template void run_batched_forward_dispatch< ck_tile::bf16_t, true, @@ -226,60 +118,6 @@ extern template void run_batched_forward_dispatch< false, 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); -extern template void run_batched_forward_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::bf16_t, - true, - false, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::bf16_t, - false, - true, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::bf16_t, - false, - true, - true, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - extern template void run_batched_forward_dispatch< ck_tile::bf16_t, true, @@ -334,60 +172,6 @@ extern template void run_batched_forward_dispatch< false, 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); -extern template void run_batched_forward_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::bf16_t, - true, - false, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::bf16_t, - false, - true, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::bf16_t, - false, - true, - true, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - extern template void run_batched_forward_dispatch< ck_tile::bf16_t, true, diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_lse_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_lse_false_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index 150aa423c6..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_lse_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_lse_false_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 0c39cd3b0d..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_lse_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_lse_false_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index e0f876c329..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index d2eae67232..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index e6882219a7..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::bf16_t, - false, - true, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 35336d77ca..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::bf16_t, - false, - true, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 27e49b03bd..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::bf16_t, - false, - true, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index 3f5bbd1603..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::bf16_t, - false, - true, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_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_lse_true_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index f516c2ae8e..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::bf16_t, - false, - true, - true, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_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_lse_true_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index b63e6ad43c..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::bf16_t, - false, - true, - true, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_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_lse_true_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index b08eb32232..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::bf16_t, - false, - true, - true, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index a9b51474ed..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::bf16_t, - false, - true, - true, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_lse_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_lse_false_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index a0428f36ca..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::fp16_t, - true, - false, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_lse_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_lse_false_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 3c683c7516..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::fp16_t, - true, - false, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_lse_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_lse_false_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index b7ea96b04c..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::fp16_t, - true, - false, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index 00d651eea2..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::fp16_t, - true, - false, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index fb558b892f..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 5b098190f6..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index b9bea58247..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index 82435aa3e6..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_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_lse_true_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index d2656b15b7..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_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_lse_true_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 3246e102fb..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_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_lse_true_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index d28ca7036d..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index 4ebbe016e5..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git 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 index 8bf2c99e79..3dff775b85 100644 --- 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 @@ -10,60 +10,6 @@ #include "hstu_attention_batched_forward_dispatch.hpp" #include "hstu_attention_params.hpp" -extern template void run_batched_forward_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::fp16_t, - true, - false, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::fp16_t, - false, - true, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::fp16_t, - false, - true, - true, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - extern template void run_batched_forward_dispatch< ck_tile::fp16_t, true, @@ -118,60 +64,6 @@ extern template void run_batched_forward_dispatch< false, 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); -extern template void run_batched_forward_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::fp16_t, - true, - false, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::fp16_t, - false, - true, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::fp16_t, - false, - true, - true, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - extern template void run_batched_forward_dispatch< ck_tile::fp16_t, true, @@ -226,60 +118,6 @@ extern template void run_batched_forward_dispatch< false, 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); -extern template void run_batched_forward_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::fp16_t, - true, - false, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::fp16_t, - false, - true, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::fp16_t, - false, - true, - true, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - extern template void run_batched_forward_dispatch< ck_tile::fp16_t, true, @@ -334,60 +172,6 @@ extern template void run_batched_forward_dispatch< false, 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); -extern template void run_batched_forward_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::fp16_t, - true, - false, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::fp16_t, - false, - true, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::fp16_t, - false, - true, - true, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - extern template void run_batched_forward_dispatch< ck_tile::fp16_t, true, diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_lse_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_lse_false_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index 1ec1dec6b9..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_lse_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_lse_false_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 7231951076..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_lse_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_lse_false_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index c1aaddcb77..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index e4544b7896..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index c450df378d..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::fp16_t, - false, - true, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index b1e01b8dfb..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::fp16_t, - false, - true, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index d94a0eaaa1..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::fp16_t, - false, - true, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index 1d4d72bf91..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::fp16_t, - false, - true, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_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_lse_true_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index 901e10cdbd..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::fp16_t, - false, - true, - true, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_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_lse_true_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index d3d7ac3934..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::fp16_t, - false, - true, - true, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_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_lse_true_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index adaf92d8e7..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::fp16_t, - false, - true, - true, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index 5feeb10b93..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_batched_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_batched_forward_dispatch< - ck_tile::fp16_t, - false, - true, - true, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index 0c6172484f..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::bf16_t, - true, - false, - false, - true, - false, - 128>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index f6840b6ea8..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::bf16_t, - true, - false, - false, - true, - false, - 256>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 3329fc7d55..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::bf16_t, - true, - false, - false, - true, - false, - 64>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index c20b780b3b..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::bf16_t, - true, - false, - false, - true, - false, - 96>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index cf01b74592..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - false, - 128>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index d4b7369b40..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - false, - 256>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 25d87823eb..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - false, - 64>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index a6b03d554b..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - false, - 96>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index ba35727fc0..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - false, - 128>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index e585679425..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - false, - 256>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 9c5d5b2b4a..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - false, - 64>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index 7c6a5ae005..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - false, - 96>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_instances_ref.hpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_instances_ref.hpp index 3ceccf6341..7b959c786d 100644 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_instances_ref.hpp +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_instances_ref.hpp @@ -10,60 +10,6 @@ #include "hstu_attention_group_forward_dispatch.hpp" #include "hstu_attention_params.hpp" -extern template void run_group_forward_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - false, - 64>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - false, - 64>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::bf16_t, - true, - false, - false, - true, - false, - 64>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::bf16_t, - false, - true, - false, - true, - false, - 64>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::bf16_t, - false, - true, - true, - true, - false, - 64>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - false, - 64>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - extern template void run_group_forward_dispatch< ck_tile::bf16_t, true, @@ -118,60 +64,6 @@ extern template void run_group_forward_dispatch< false, 64>(HstuAttentionGroupFwdParams& param, hipStream_t stream); -extern template void run_group_forward_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - false, - 96>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - false, - 96>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::bf16_t, - true, - false, - false, - true, - false, - 96>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::bf16_t, - false, - true, - false, - true, - false, - 96>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::bf16_t, - false, - true, - true, - true, - false, - 96>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - false, - 96>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - extern template void run_group_forward_dispatch< ck_tile::bf16_t, true, @@ -226,60 +118,6 @@ extern template void run_group_forward_dispatch< false, 96>(HstuAttentionGroupFwdParams& param, hipStream_t stream); -extern template void run_group_forward_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - false, - 128>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - false, - 128>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::bf16_t, - true, - false, - false, - true, - false, - 128>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::bf16_t, - false, - true, - false, - true, - false, - 128>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::bf16_t, - false, - true, - true, - true, - false, - 128>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - false, - 128>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - extern template void run_group_forward_dispatch< ck_tile::bf16_t, true, @@ -334,60 +172,6 @@ extern template void run_group_forward_dispatch< false, 128>(HstuAttentionGroupFwdParams& param, hipStream_t stream); -extern template void run_group_forward_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - false, - 256>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - false, - 256>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::bf16_t, - true, - false, - false, - true, - false, - 256>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::bf16_t, - false, - true, - false, - true, - false, - 256>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::bf16_t, - false, - true, - true, - true, - false, - 256>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - false, - 256>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - extern template void run_group_forward_dispatch< ck_tile::bf16_t, true, diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index 4e779dcfe7..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - false, - 128>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 096db3e82b..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - false, - 256>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 03207625c7..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - false, - 64>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index e626a36e22..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - false, - 96>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index 7e1aa9de58..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::bf16_t, - false, - true, - false, - true, - false, - 128>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 72dd1bef1c..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::bf16_t, - false, - true, - false, - true, - false, - 256>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index afcb49ebf1..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::bf16_t, - false, - true, - false, - true, - false, - 64>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index 6ef12d51a0..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::bf16_t, - false, - true, - false, - true, - false, - 96>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index 29f6c79249..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::bf16_t, - false, - true, - true, - true, - false, - 128>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 182c113068..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::bf16_t, - false, - true, - true, - true, - false, - 256>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 1f3b2f18f8..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::bf16_t, - false, - true, - true, - true, - false, - 64>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index 9355064c7c..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::bf16_t, - false, - true, - true, - true, - false, - 96>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index 4fa0b9ac83..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::fp16_t, - true, - false, - false, - true, - false, - 128>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index a3b759555b..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::fp16_t, - true, - false, - false, - true, - false, - 256>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 2df2a4b29c..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::fp16_t, - true, - false, - false, - true, - false, - 64>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index 11924879b0..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::fp16_t, - true, - false, - false, - true, - false, - 96>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index 8663b42c0a..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - false, - 128>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 901b92cdf4..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - false, - 256>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index bf3db24504..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - false, - 64>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index 85c4fb3b81..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - false, - 96>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index 3482baf46c..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - false, - 128>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index b9445a6d4b..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - false, - 256>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 7eb54ef8a8..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - false, - 64>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index ef34b51f18..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - false, - 96>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_instances_ref.hpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_instances_ref.hpp index 8a1dfd6131..b07e91eca1 100644 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_instances_ref.hpp +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_instances_ref.hpp @@ -10,60 +10,6 @@ #include "hstu_attention_group_forward_dispatch.hpp" #include "hstu_attention_params.hpp" -extern template void run_group_forward_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - false, - 64>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - false, - 64>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::fp16_t, - true, - false, - false, - true, - false, - 64>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::fp16_t, - false, - true, - false, - true, - false, - 64>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::fp16_t, - false, - true, - true, - true, - false, - 64>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - false, - 64>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - extern template void run_group_forward_dispatch< ck_tile::fp16_t, true, @@ -118,60 +64,6 @@ extern template void run_group_forward_dispatch< false, 64>(HstuAttentionGroupFwdParams& param, hipStream_t stream); -extern template void run_group_forward_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - false, - 96>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - false, - 96>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::fp16_t, - true, - false, - false, - true, - false, - 96>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::fp16_t, - false, - true, - false, - true, - false, - 96>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::fp16_t, - false, - true, - true, - true, - false, - 96>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - false, - 96>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - extern template void run_group_forward_dispatch< ck_tile::fp16_t, true, @@ -226,60 +118,6 @@ extern template void run_group_forward_dispatch< false, 96>(HstuAttentionGroupFwdParams& param, hipStream_t stream); -extern template void run_group_forward_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - false, - 128>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - false, - 128>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::fp16_t, - true, - false, - false, - true, - false, - 128>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::fp16_t, - false, - true, - false, - true, - false, - 128>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::fp16_t, - false, - true, - true, - true, - false, - 128>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - false, - 128>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - extern template void run_group_forward_dispatch< ck_tile::fp16_t, true, @@ -334,60 +172,6 @@ extern template void run_group_forward_dispatch< false, 128>(HstuAttentionGroupFwdParams& param, hipStream_t stream); -extern template void run_group_forward_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - false, - 256>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - false, - 256>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::fp16_t, - true, - false, - false, - true, - false, - 256>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::fp16_t, - false, - true, - false, - true, - false, - 256>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::fp16_t, - false, - true, - true, - true, - false, - 256>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - -extern template void run_group_forward_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - false, - 256>(HstuAttentionGroupFwdParams& param, hipStream_t stream); - extern template void run_group_forward_dispatch< ck_tile::fp16_t, true, diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index ff124bf4c3..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - false, - 128>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index c4d067bb07..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - false, - 256>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index ffbbcf2106..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - false, - 64>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index 43ba4fad88..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - false, - 96>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index ab0c4fdac6..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::fp16_t, - false, - true, - false, - true, - false, - 128>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index cc4e39bde1..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::fp16_t, - false, - true, - false, - true, - false, - 256>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 85b37925f4..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::fp16_t, - false, - true, - false, - true, - false, - 64>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index ccd4e36e1c..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::fp16_t, - false, - true, - false, - true, - false, - 96>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index 572850e87f..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::fp16_t, - false, - true, - true, - true, - false, - 128>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 59f0d24f0d..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::fp16_t, - false, - true, - true, - true, - false, - 256>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index bcef68b995..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::fp16_t, - false, - true, - true, - true, - false, - 64>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index 325675b532..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_group_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_group_forward_dispatch< - ck_tile::fp16_t, - false, - true, - true, - true, - false, - 96>(HstuAttentionGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index 9d49229841..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - true, - false, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 824b4e0585..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - true, - false, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index ae63583e54..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - true, - false, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index 9939d0e268..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - true, - false, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index 757d94bdc2..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index d449376faf..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 1a3e0f3d0d..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index d333e05ab8..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index 0bffc4baef..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 74e1b2bde2..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index ef9a4524ad..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index eb74e0c343..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_instances_ref.hpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_instances_ref.hpp index b66ff7f2cd..67155df332 100644 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_instances_ref.hpp +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_instances_ref.hpp @@ -10,60 +10,6 @@ #include "hstu_attention_jagged_forward_dispatch.hpp" #include "hstu_attention_params.hpp" -extern template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - true, - false, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - false, - true, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - false, - true, - true, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - extern template void run_jagged_forward_dispatch< ck_tile::bf16_t, true, @@ -118,60 +64,6 @@ extern template void run_jagged_forward_dispatch< false, 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); -extern template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - true, - false, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - false, - true, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - false, - true, - true, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - extern template void run_jagged_forward_dispatch< ck_tile::bf16_t, true, @@ -226,60 +118,6 @@ extern template void run_jagged_forward_dispatch< false, 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); -extern template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - true, - false, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - false, - true, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - false, - true, - true, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - extern template void run_jagged_forward_dispatch< ck_tile::bf16_t, true, @@ -334,60 +172,6 @@ extern template void run_jagged_forward_dispatch< false, 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); -extern template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - true, - false, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - false, - true, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - false, - true, - true, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - extern template void run_jagged_forward_dispatch< ck_tile::bf16_t, true, diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index 9e27c3c524..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 24497a3306..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 11f6e36444..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index bfcc992625..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index ff90df43e3..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - false, - true, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 4f2089089c..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - false, - true, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 7414c05c9f..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - false, - true, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index 584d57797f..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - false, - true, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index 3753ad3383..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - false, - true, - true, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index db8ab2e268..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - false, - true, - true, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 9e09a70f4e..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - false, - true, - true, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index b4133fa129..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::bf16_t, - false, - true, - true, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index 6338de98ce..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - true, - false, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 5122bcd134..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - true, - false, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index c6ec65e72f..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - true, - false, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index 8e8553c710..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - true, - false, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index c18fbb9701..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 7fca09c2b3..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 9c122b57d2..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index c3f51a6b67..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index c9c100e856..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index a920792ed3..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 399a4776a4..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index 04efc0cac4..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_instances_ref.hpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_instances_ref.hpp index 5b9799744b..6275cf2bc4 100644 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_instances_ref.hpp +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_instances_ref.hpp @@ -10,60 +10,6 @@ #include "hstu_attention_jagged_forward_dispatch.hpp" #include "hstu_attention_params.hpp" -extern template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - true, - false, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - false, - true, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - false, - true, - true, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - extern template void run_jagged_forward_dispatch< ck_tile::fp16_t, true, @@ -118,60 +64,6 @@ extern template void run_jagged_forward_dispatch< false, 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); -extern template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - true, - false, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - false, - true, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - false, - true, - true, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - extern template void run_jagged_forward_dispatch< ck_tile::fp16_t, true, @@ -226,60 +118,6 @@ extern template void run_jagged_forward_dispatch< false, 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); -extern template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - true, - false, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - false, - true, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - false, - true, - true, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - extern template void run_jagged_forward_dispatch< ck_tile::fp16_t, true, @@ -334,60 +172,6 @@ extern template void run_jagged_forward_dispatch< false, 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); -extern template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - true, - false, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - false, - true, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - false, - true, - true, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); - extern template void run_jagged_forward_dispatch< ck_tile::fp16_t, true, diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index 9937ef1232..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index ae71e4fd52..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 6e8b7b489d..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index 5f17b9d109..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_false_lse_false_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index 0b22e77934..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - false, - true, - false, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 20ffdd9ded..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - false, - true, - false, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 970eb2a7d8..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - false, - true, - false, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index 25e8ad33d7..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_false_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - false, - true, - false, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp deleted file mode 100644 index a6e6e6cca3..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_128.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - false, - true, - true, - true, - false, - 128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 84269bcc32..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - false, - true, - true, - true, - false, - 256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index b32c0425e3..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - false, - true, - true, - true, - false, - 64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp deleted file mode 100644 index b8cc09b0ef..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_softmax_true_lse_true_has_bias_no_dropout_maxk_96.cpp +++ /dev/null @@ -1,20 +0,0 @@ - -// 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 -#include "hstu_attention_jagged_forward_dispatch.hpp" -#include "hstu_attention_params.hpp" - -template void run_jagged_forward_dispatch< - ck_tile::fp16_t, - false, - true, - true, - true, - false, - 96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);