diff --git a/example/ck_tile/18_hstu_attention/generate_instances.py b/example/ck_tile/18_hstu_attention/generate_instances.py index 0f6f6191fd..487499faf0 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..08ef2acc19 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 @@ -8,21 +8,11 @@ #define HDIM_SWITCH(HDIM_1, HDIM_2, CONST_NAME, ...) \ [&] { \ - if(HDIM_1 <= 64 && HDIM_2 <= 64) \ - { \ - constexpr ck_tile::index_t CONST_NAME = 64; \ - __VA_ARGS__(); \ - } \ - else if(HDIM_1 <= 128 && HDIM_2 <= 128) \ + if(HDIM_1 <= 128 && HDIM_2 <= 128) \ { \ constexpr ck_tile::index_t CONST_NAME = 128; \ __VA_ARGS__(); \ } \ - else if(HDIM_1 <= 256 && HDIM_2 <= 256) \ - { \ - constexpr ck_tile::index_t CONST_NAME = 256; \ - __VA_ARGS__(); \ - } \ else \ { \ throw std::runtime_error("Head-dim sizes not supported!"); \ 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_has_causal_no_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index 9bbdd62a9f..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_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, - true, - false, - 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_no_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index 4a65433223..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_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, - true, - false, - 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_no_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index bee6cb5b7b..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_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, - true, - false, - 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_no_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index e0ece6745a..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_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, - true, - false, - 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_no_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index a0fa6eaf31..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_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, - true, - false, - 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_no_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index 49b1192f4f..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_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, - true, - false, - 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_no_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index a5da3f40ab..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_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, - true, - false, - 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_no_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 2ba44c444e..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_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, - true, - false, - 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_has_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index 450a70774d..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_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, - false, - true, - 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_has_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index bdc14f4caa..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_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, - false, - true, - 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_has_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 2bb09b84ca..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_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, - false, - true, - 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_has_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 24db78fdab..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_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, - false, - true, - 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_has_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index 7788e7f89b..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_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, - false, - true, - 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_has_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index 6227be13ac..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_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, - false, - true, - 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_has_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 566da83708..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_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, - false, - true, - 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_has_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index f39d11a2aa..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_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, - false, - true, - false, - 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_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_has_causal_no_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index 693fe022f6..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_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, - true, - false, - 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_no_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index 73cb991466..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_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, - true, - false, - 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_no_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 70d9a31ae0..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_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, - true, - false, - 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_no_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 51acf1b238..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_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, - true, - false, - 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_no_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index 295001b585..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_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, - true, - false, - 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_no_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index 3658825391..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_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, - true, - false, - 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_no_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 4078670fd9..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_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, - true, - false, - 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_no_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 314f577c91..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_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, - true, - false, - 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_has_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index ab436fc646..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_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, - false, - true, - 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_has_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index d2cc4b579e..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_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, - false, - true, - 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_has_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 71bb7a409d..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_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, - false, - true, - 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_has_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 0f209d16e6..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_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, - false, - true, - 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_has_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index cbefee339d..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_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, - false, - true, - 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_has_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index 84fe9cf731..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_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, - false, - true, - 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_has_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index d0f8a1131a..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_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, - false, - true, - 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_has_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index a1eb2ed6f6..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_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, - false, - true, - false, - 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_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_has_causal_no_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index 58bc729e42..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_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, - true, - false, - 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_no_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index 6e8ba77772..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_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, - true, - false, - 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_no_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index f717ee4563..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_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, - true, - false, - 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_no_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 258b1a576e..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_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, - true, - false, - 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_no_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index d9b4631acd..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_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, - true, - false, - 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_no_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index 7796f67298..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_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, - true, - false, - 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_no_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index a442eda92b..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_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, - true, - false, - 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_no_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index d79f8ffacd..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_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, - true, - false, - 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_has_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index cbccdcdf0a..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_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, - false, - true, - 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_has_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index ea329e073f..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_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, - false, - true, - 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_has_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 614eae2111..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_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, - false, - true, - 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_has_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 428e08cc13..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_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, - false, - true, - 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_has_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index aa35336a75..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_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, - false, - true, - 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_has_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index 3070211e3f..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_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, - false, - true, - 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_has_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 44abb0207f..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_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, - false, - true, - 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_has_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index b08d8d7012..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_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, - false, - true, - false, - 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_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_has_causal_no_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_has_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index 68af8bdfaf..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_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, - true, - false, - 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_no_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_has_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index 0bedd49441..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_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, - true, - false, - 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_no_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 3679db4777..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_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, - true, - false, - 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_no_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index dd9ad9c2bf..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_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, - true, - false, - 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_no_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_no_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index 33d24204a9..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_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, - true, - false, - 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_no_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_no_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index 12f1013041..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_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, - true, - false, - 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_no_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_no_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 4693f6ea05..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_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, - true, - false, - 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_no_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_no_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index c4427d8e6e..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_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, - true, - false, - 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_has_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_has_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index 00be5745df..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_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, - false, - true, - 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_has_local_has_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_has_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index 2cb1cd7e83..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_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, - false, - true, - 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_has_local_has_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_has_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index 25b0dc246e..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_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, - false, - true, - 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_has_local_has_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_has_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index a4b8bb18e8..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_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, - false, - true, - 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_has_local_no_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_no_bias_has_dropout_maxk_256.cpp deleted file mode 100644 index af606140ed..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_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, - false, - true, - 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_has_local_no_bias_has_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_no_bias_has_dropout_maxk_64.cpp deleted file mode 100644 index 4b05b4e9b7..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_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, - false, - true, - 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_has_local_no_bias_no_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_no_bias_no_dropout_maxk_256.cpp deleted file mode 100644 index b14466864c..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_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, - false, - true, - 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_has_local_no_bias_no_dropout_maxk_64.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_no_bias_no_dropout_maxk_64.cpp deleted file mode 100644 index 04e311ca20..0000000000 --- a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_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, - false, - true, - false, - 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_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);