From 2972de4c88a3e11e63b2ec3b0014ce8aba911ad6 Mon Sep 17 00:00:00 2001 From: Qianfeng Zhang Date: Wed, 30 Apr 2025 01:54:43 +0000 Subject: [PATCH] Temporarily close the instance for hdim64 and hdim256 to save compiling time --- .../18_hstu_attention/generate_instances.py | 2 +- .../hstu_attention_hdim_switch.hpp | 4 +- ...as_local_has_bias_has_dropout_maxk_256.cpp | 22 --- ...has_local_has_bias_has_dropout_maxk_64.cpp | 22 --- ...has_local_has_bias_no_dropout_maxk_256.cpp | 22 --- ..._has_local_has_bias_no_dropout_maxk_64.cpp | 22 --- ...has_local_no_bias_has_dropout_maxk_256.cpp | 22 --- ..._has_local_no_bias_has_dropout_maxk_64.cpp | 22 --- ..._has_local_no_bias_no_dropout_maxk_256.cpp | 22 --- ...l_has_local_no_bias_no_dropout_maxk_64.cpp | 22 --- ...ion_batched_forward_bf16_instances_ref.hpp | 128 ------------------ ...no_local_has_bias_has_dropout_maxk_256.cpp | 22 --- ..._no_local_has_bias_has_dropout_maxk_64.cpp | 22 --- ..._no_local_has_bias_no_dropout_maxk_256.cpp | 22 --- ...l_no_local_has_bias_no_dropout_maxk_64.cpp | 22 --- ..._no_local_no_bias_has_dropout_maxk_256.cpp | 22 --- ...l_no_local_no_bias_has_dropout_maxk_64.cpp | 22 --- ...l_no_local_no_bias_no_dropout_maxk_256.cpp | 22 --- ...al_no_local_no_bias_no_dropout_maxk_64.cpp | 22 --- ...as_local_has_bias_has_dropout_maxk_256.cpp | 22 --- ...has_local_has_bias_has_dropout_maxk_64.cpp | 22 --- ...has_local_has_bias_no_dropout_maxk_256.cpp | 22 --- ..._has_local_has_bias_no_dropout_maxk_64.cpp | 22 --- ...has_local_no_bias_has_dropout_maxk_256.cpp | 22 --- ..._has_local_no_bias_has_dropout_maxk_64.cpp | 22 --- ..._has_local_no_bias_no_dropout_maxk_256.cpp | 22 --- ...l_has_local_no_bias_no_dropout_maxk_64.cpp | 22 --- ...ion_batched_forward_fp16_instances_ref.hpp | 128 ------------------ ...no_local_has_bias_has_dropout_maxk_256.cpp | 22 --- ..._no_local_has_bias_has_dropout_maxk_64.cpp | 22 --- ..._no_local_has_bias_no_dropout_maxk_256.cpp | 22 --- ...l_no_local_has_bias_no_dropout_maxk_64.cpp | 22 --- ..._no_local_no_bias_has_dropout_maxk_256.cpp | 22 --- ...l_no_local_no_bias_has_dropout_maxk_64.cpp | 22 --- ...l_no_local_no_bias_no_dropout_maxk_256.cpp | 22 --- ...al_no_local_no_bias_no_dropout_maxk_64.cpp | 22 --- ...as_local_has_bias_has_dropout_maxk_256.cpp | 22 --- ...has_local_has_bias_has_dropout_maxk_64.cpp | 22 --- ...has_local_has_bias_no_dropout_maxk_256.cpp | 22 --- ..._has_local_has_bias_no_dropout_maxk_64.cpp | 22 --- ...has_local_no_bias_has_dropout_maxk_256.cpp | 22 --- ..._has_local_no_bias_has_dropout_maxk_64.cpp | 22 --- ..._has_local_no_bias_no_dropout_maxk_256.cpp | 22 --- ...l_has_local_no_bias_no_dropout_maxk_64.cpp | 22 --- ...tion_jagged_forward_bf16_instances_ref.hpp | 128 ------------------ ...no_local_has_bias_has_dropout_maxk_256.cpp | 22 --- ..._no_local_has_bias_has_dropout_maxk_64.cpp | 22 --- ..._no_local_has_bias_no_dropout_maxk_256.cpp | 22 --- ...l_no_local_has_bias_no_dropout_maxk_64.cpp | 22 --- ..._no_local_no_bias_has_dropout_maxk_256.cpp | 22 --- ...l_no_local_no_bias_has_dropout_maxk_64.cpp | 22 --- ...l_no_local_no_bias_no_dropout_maxk_256.cpp | 22 --- ...al_no_local_no_bias_no_dropout_maxk_64.cpp | 22 --- ...as_local_has_bias_has_dropout_maxk_256.cpp | 22 --- ...has_local_has_bias_has_dropout_maxk_64.cpp | 22 --- ...has_local_has_bias_no_dropout_maxk_256.cpp | 22 --- ..._has_local_has_bias_no_dropout_maxk_64.cpp | 22 --- ...has_local_no_bias_has_dropout_maxk_256.cpp | 22 --- ..._has_local_no_bias_has_dropout_maxk_64.cpp | 22 --- ..._has_local_no_bias_no_dropout_maxk_256.cpp | 22 --- ...l_has_local_no_bias_no_dropout_maxk_64.cpp | 22 --- ...tion_jagged_forward_fp16_instances_ref.hpp | 128 ------------------ ...no_local_has_bias_has_dropout_maxk_256.cpp | 22 --- ..._no_local_has_bias_has_dropout_maxk_64.cpp | 22 --- ..._no_local_has_bias_no_dropout_maxk_256.cpp | 22 --- ...l_no_local_has_bias_no_dropout_maxk_64.cpp | 22 --- ..._no_local_no_bias_has_dropout_maxk_256.cpp | 22 --- ...l_no_local_no_bias_has_dropout_maxk_64.cpp | 22 --- ...l_no_local_no_bias_no_dropout_maxk_256.cpp | 22 --- ...al_no_local_no_bias_no_dropout_maxk_64.cpp | 22 --- 70 files changed, 3 insertions(+), 1923 deletions(-) delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp delete mode 100644 example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp diff --git a/example/ck_tile/18_hstu_attention/generate_instances.py b/example/ck_tile/18_hstu_attention/generate_instances.py index 3427d1adf8..e5ea9b2ce1 100644 --- a/example/ck_tile/18_hstu_attention/generate_instances.py +++ b/example/ck_tile/18_hstu_attention/generate_instances.py @@ -161,7 +161,7 @@ def create_forward_instances_ref(instance_dir: Path, headdims: List) -> None: file.write(forward_instance) if __name__ == "__main__": - headdims_fwd = [64, 128, 256] + headdims_fwd = [128] this_dir = os.path.dirname(__file__) output_dir = Path(this_dir) / "instances" diff --git a/example/ck_tile/18_hstu_attention/hstu_attention_hdim_switch.hpp b/example/ck_tile/18_hstu_attention/hstu_attention_hdim_switch.hpp index ab8cbe31ed..43be47795c 100644 --- a/example/ck_tile/18_hstu_attention/hstu_attention_hdim_switch.hpp +++ b/example/ck_tile/18_hstu_attention/hstu_attention_hdim_switch.hpp @@ -10,7 +10,7 @@ [&] { \ if(HDIM_1 <= 64 && HDIM_2 <= 64) \ { \ - constexpr ck_tile::index_t CONST_NAME = 64; \ + constexpr ck_tile::index_t CONST_NAME = 128; \ __VA_ARGS__(); \ } \ else if(HDIM_1 <= 128 && HDIM_2 <= 128) \ @@ -20,7 +20,7 @@ } \ else if(HDIM_1 <= 256 && HDIM_2 <= 256) \ { \ - constexpr ck_tile::index_t CONST_NAME = 256; \ + constexpr ck_tile::index_t CONST_NAME = 128; \ __VA_ARGS__(); \ } \ else \ diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index 1240e37c28..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index dc61eb8ce7..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index fecc66f9ef..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - true, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index ee6bc64f14..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - true, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index e567c295f4..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index 4a9fdb3052..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index a75ff8e76a..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - false, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index a4f6adc2eb..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - false, - false, - 64>(HstuAttentionFwdParams& 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 c4a377f9ed..c869807db9 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 @@ -13,70 +13,6 @@ #include #include "hstu_attention_batched_forward_dispatch.hpp" -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - true, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - true, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - true, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - false, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - false, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - extern template void run_batched_forward_causal_local_bias_dropout_dispatch< ck_tile::bf16_t, true, @@ -140,67 +76,3 @@ extern template void run_batched_forward_causal_local_bias_dropout_dispatch< false, false, 128>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - true, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - true, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - true, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - false, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - false, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index f723f0fb38..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - true, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index 09be671b6c..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - true, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 6b2de8cc19..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - true, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index fe1e8b3509..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - true, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index fd41687481..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index 89160eb715..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 1f4508cd6b..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - false, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index dd0c7978f2..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - false, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index 2bba03d1e9..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index a10e926a29..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index dda7dc07aa..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - true, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index ce48b9a401..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - true, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index 6e700fce4b..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index 5c0206156c..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index f2f1bdc3cd..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - false, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index cd9979a4ee..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - false, - false, - 64>(HstuAttentionFwdParams& 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 5c350254dc..38fd44a3e2 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 @@ -13,70 +13,6 @@ #include #include "hstu_attention_batched_forward_dispatch.hpp" -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - true, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - true, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - true, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - false, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - false, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - extern template void run_batched_forward_causal_local_bias_dropout_dispatch< ck_tile::fp16_t, true, @@ -140,67 +76,3 @@ extern template void run_batched_forward_causal_local_bias_dropout_dispatch< false, false, 128>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - true, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - true, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - true, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - false, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - false, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index d399cf6dbb..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - true, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index 91875c4deb..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - true, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 653c378852..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - true, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index c926f2e9b7..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - true, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index 2d60c18f02..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index bbf5f6817e..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 411afb0ec1..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - false, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 49bb0b3ce4..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_batched_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - false, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index c1e80c9666..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index 795c8cc4ed..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 0f3ee275fb..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - true, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index c875587af3..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - true, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index 8e73023e0c..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index 71050ae3d7..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index bf40833a7f..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - false, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index cef802a33e..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - false, - false, - 64>(HstuAttentionFwdParams& 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 0bdbdd91fc..44a80d98a8 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 @@ -13,70 +13,6 @@ #include #include "hstu_attention_jagged_forward_dispatch.hpp" -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - true, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - true, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - true, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - false, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - false, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< ck_tile::bf16_t, true, @@ -140,67 +76,3 @@ extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< false, false, 128>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - true, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - true, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - true, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - true, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - false, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - true, - true, - false, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - false, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index 4f330cd73b..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - true, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index 22ca2c84ad..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - true, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 6a06c0e0d3..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - true, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 3bbbe402d2..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - true, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index 8337a1559f..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index 3c09e8416f..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - false, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index ab45759547..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - false, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 2843a7a1c5..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::bf16_t, - false, - false, - false, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index 10b05c8a77..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index 83d43f43ef..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 5093f4742f..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - true, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 60440c4916..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - true, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index 1d55f69805..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index ea8e8895e9..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 88464a5075..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - false, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 5be82db604..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - false, - false, - 64>(HstuAttentionFwdParams& 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 c4ae5f851b..1eac6f27e7 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 @@ -13,70 +13,6 @@ #include #include "hstu_attention_jagged_forward_dispatch.hpp" -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - true, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - true, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - true, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - false, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - false, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); - extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< ck_tile::fp16_t, true, @@ -140,67 +76,3 @@ extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< false, false, 128>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - true, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - true, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - true, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - true, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - false, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - true, - true, - false, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); - -extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - false, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index 23de813bbe..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - true, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index 8aead96641..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - true, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index a36b542f6d..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - true, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 59d20c38e7..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - true, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index c78ee824fb..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index 195fa6cc0a..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - false, - true, - 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 0a7a51034d..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - false, - false, - 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 31fbe9459c..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp +++ /dev/null @@ -1,22 +0,0 @@ - -/* - Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. - * - * This source code is licensed under the BSD-style license found in the - * LICENSE file in the root directory of this source tree. - * - * 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" - -template void run_jagged_forward_causal_local_bias_dropout_dispatch< - ck_tile::fp16_t, - false, - false, - false, - false, - 64>(HstuAttentionFwdParams& param, hipStream_t stream);