diff --git a/example/ck_tile/18_hstu_attention/generate_instances.py b/example/ck_tile/18_hstu_attention/generate_instances.py index e5ea9b2ce1..3427d1adf8 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 = [128] + headdims_fwd = [64, 128, 256] 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 43be47795c..ab8cbe31ed 100644 --- a/example/ck_tile/18_hstu_attention/hstu_attention_hdim_switch.hpp +++ b/example/ck_tile/18_hstu_attention/hstu_attention_hdim_switch.hpp @@ -10,7 +10,7 @@ [&] { \ if(HDIM_1 <= 64 && HDIM_2 <= 64) \ { \ - constexpr ck_tile::index_t CONST_NAME = 128; \ + constexpr ck_tile::index_t CONST_NAME = 64; \ __VA_ARGS__(); \ } \ else if(HDIM_1 <= 128 && HDIM_2 <= 128) \ @@ -20,7 +20,7 @@ } \ else if(HDIM_1 <= 256 && HDIM_2 <= 256) \ { \ - constexpr ck_tile::index_t CONST_NAME = 128; \ + constexpr ck_tile::index_t CONST_NAME = 256; \ __VA_ARGS__(); \ } \ else \ diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..1240e37c28 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..dc61eb8ce7 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..fecc66f9ef --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..ee6bc64f14 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..e567c295f4 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..4a9fdb3052 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..a75ff8e76a --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..a4f6adc2eb --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_instances_ref.hpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_instances_ref.hpp index c869807db9..c4a377f9ed 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,6 +13,70 @@ #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, @@ -76,3 +140,67 @@ extern template void run_batched_forward_causal_local_bias_dropout_dispatch< false, false, 128>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..f723f0fb38 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..09be671b6c --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..6b2de8cc19 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..fe1e8b3509 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..fd41687481 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..89160eb715 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..1f4508cd6b --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..dd0c7978f2 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..2bba03d1e9 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..a10e926a29 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..dda7dc07aa --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..ce48b9a401 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..6e700fce4b --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..5c0206156c --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..f2f1bdc3cd --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..cd9979a4ee --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_batched_forward_dispatch.hpp" + +template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_instances_ref.hpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_instances_ref.hpp index 38fd44a3e2..5c350254dc 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,6 +13,70 @@ #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, @@ -76,3 +140,67 @@ extern template void run_batched_forward_causal_local_bias_dropout_dispatch< false, false, 128>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_batched_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..d399cf6dbb --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..91875c4deb --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..653c378852 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..c926f2e9b7 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..2d60c18f02 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..bbf5f6817e --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..411afb0ec1 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..49bb0b3ce4 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..c1e80c9666 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..795c8cc4ed --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..0f3ee275fb --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..c875587af3 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..8e73023e0c --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..71050ae3d7 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..bf40833a7f --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..cef802a33e --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_instances_ref.hpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_instances_ref.hpp index 44a80d98a8..0bdbdd91fc 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,6 +13,70 @@ #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, @@ -76,3 +140,67 @@ extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< false, false, 128>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + true, + true, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::bf16_t, + false, + false, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..4f330cd73b --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..22ca2c84ad --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..6a06c0e0d3 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..3bbbe402d2 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..8337a1559f --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..3c09e8416f --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..ab45759547 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..2843a7a1c5 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..10b05c8a77 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..83d43f43ef --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..5093f4742f --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..60440c4916 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..1d55f69805 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..ea8e8895e9 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..88464a5075 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..5be82db604 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved. + * + * This source code is licensed under the BSD-style license found in the + * LICENSE file in the root directory of this source tree. + * + * The file is automatically generated, don't modify! + * See the generator script + * `composable_kernel/example/ck_tile/18_hstu_attention/generate_instances.py` + */ + +#include +#include "hstu_attention_jagged_forward_dispatch.hpp" + +template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_instances_ref.hpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_instances_ref.hpp index 1eac6f27e7..c4ae5f851b 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,6 +13,70 @@ #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, @@ -76,3 +140,67 @@ extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< false, false, 128>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + true, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + true, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + false, + true, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + true, + true, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); + +extern template void run_jagged_forward_causal_local_bias_dropout_dispatch< + ck_tile::fp16_t, + false, + false, + false, + false, + 256>(HstuAttentionFwdParams& param, hipStream_t stream); diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp new file mode 100644 index 0000000000..23de813bbe --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..8aead96641 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..a36b542f6d --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..59d20c38e7 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..c78ee824fb --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..195fa6cc0a --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..0a7a51034d --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp @@ -0,0 +1,22 @@ + +/* + 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 new file mode 100644 index 0000000000..31fbe9459c --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp @@ -0,0 +1,22 @@ + +/* + 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);