diff --git a/example/ck_tile/18_hstu_attention/generate_instances.py b/example/ck_tile/18_hstu_attention/generate_instances.py index 3427d1adf8..0f6f6191fd 100644 --- a/example/ck_tile/18_hstu_attention/generate_instances.py +++ b/example/ck_tile/18_hstu_attention/generate_instances.py @@ -88,7 +88,7 @@ MODE_NAME_MAP = { def create_forward_instances(instance_dir: Path, headdims: List) -> None: for mode in ["batched", "jagged"]: for dtype in ["fp16", "bf16"]: - for has_causal, has_local in zip([True, False],[True, False]): + for has_causal, has_local in ([True, True], [True, False], [False, True], [False, False]): for has_bias in [True, False]: for has_dropout in [True, False]: for max_k in headdims: @@ -112,7 +112,7 @@ def create_forward_instances(instance_dir: Path, headdims: List) -> None: mode=mode, dtype=TYPE_CTYPE_MAP[dtype], has_causal=BOOL_MAP[has_causal], - has_local=BOOL_MAP[has_causal], + has_local=BOOL_MAP[has_local], has_bias=BOOL_MAP[has_bias], has_dropout=BOOL_MAP[has_dropout], max_k=max_k, diff --git a/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_128.cpp new file mode 100644 index 0000000000..8924b97631 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_128.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, + false, + true, + true, + 128>(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 new file mode 100644 index 0000000000..9bbdd62a9f --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_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, + 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 new file mode 100644 index 0000000000..4a65433223 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_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, + 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_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_128.cpp new file mode 100644 index 0000000000..1473328181 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_128.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, + false, + true, + false, + 128>(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 new file mode 100644 index 0000000000..bee6cb5b7b --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_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, + 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 new file mode 100644 index 0000000000..e0ece6745a --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_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, + 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_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_128.cpp new file mode 100644 index 0000000000..52f27706bc --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_128.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, + false, + false, + true, + 128>(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 new file mode 100644 index 0000000000..a0fa6eaf31 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_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, + 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 new file mode 100644 index 0000000000..49b1192f4f --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_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, + 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_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_128.cpp new file mode 100644 index 0000000000..7581915e4b --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_128.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, + false, + false, + false, + 128>(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 new file mode 100644 index 0000000000..a5da3f40ab --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_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, + 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 new file mode 100644 index 0000000000..2ba44c444e --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_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, + 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_no_causal_has_local_has_bias_has_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_128.cpp new file mode 100644 index 0000000000..c8fcc94019 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_128.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, + true, + true, + true, + 128>(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 new file mode 100644 index 0000000000..450a70774d --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_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, + 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 new file mode 100644 index 0000000000..bdc14f4caa --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_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, + 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_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_128.cpp new file mode 100644 index 0000000000..c9820c4df6 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_128.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, + true, + true, + false, + 128>(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 new file mode 100644 index 0000000000..2bb09b84ca --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_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, + 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 new file mode 100644 index 0000000000..24db78fdab --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_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, + 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_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_128.cpp new file mode 100644 index 0000000000..69c2ec0d5a --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_128.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, + true, + false, + true, + 128>(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 new file mode 100644 index 0000000000..7788e7f89b --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_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, + 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 new file mode 100644 index 0000000000..6227be13ac --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_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, + 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_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_128.cpp new file mode 100644 index 0000000000..fcd5c78b0a --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_128.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, + true, + false, + false, + 128>(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 new file mode 100644 index 0000000000..566da83708 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_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, + 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 new file mode 100644 index 0000000000..f39d11a2aa --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_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, + 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_has_causal_no_local_has_bias_has_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_has_dropout_maxk_128.cpp new file mode 100644 index 0000000000..c2c12dac01 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_has_dropout_maxk_128.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, + false, + true, + true, + 128>(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 new file mode 100644 index 0000000000..693fe022f6 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_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, + 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 new file mode 100644 index 0000000000..73cb991466 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_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, + 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_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_no_dropout_maxk_128.cpp new file mode 100644 index 0000000000..120b531518 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_no_dropout_maxk_128.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, + false, + true, + false, + 128>(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 new file mode 100644 index 0000000000..70d9a31ae0 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_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, + 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 new file mode 100644 index 0000000000..51acf1b238 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_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, + 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_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_has_dropout_maxk_128.cpp new file mode 100644 index 0000000000..42495b7675 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_has_dropout_maxk_128.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, + false, + false, + true, + 128>(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 new file mode 100644 index 0000000000..295001b585 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_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, + 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 new file mode 100644 index 0000000000..3658825391 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_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, + 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_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_no_dropout_maxk_128.cpp new file mode 100644 index 0000000000..b1b8134ea6 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_no_dropout_maxk_128.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, + false, + false, + false, + 128>(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 new file mode 100644 index 0000000000..4078670fd9 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_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, + 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 new file mode 100644 index 0000000000..314f577c91 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_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, + 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_no_causal_has_local_has_bias_has_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_has_dropout_maxk_128.cpp new file mode 100644 index 0000000000..2e4817a09c --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_has_dropout_maxk_128.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, + true, + true, + true, + 128>(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 new file mode 100644 index 0000000000..ab436fc646 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_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, + 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 new file mode 100644 index 0000000000..d2cc4b579e --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_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, + 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_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_no_dropout_maxk_128.cpp new file mode 100644 index 0000000000..40a27d584c --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_no_dropout_maxk_128.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, + true, + true, + false, + 128>(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 new file mode 100644 index 0000000000..71bb7a409d --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_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, + 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 new file mode 100644 index 0000000000..0f209d16e6 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_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, + 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_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_has_dropout_maxk_128.cpp new file mode 100644 index 0000000000..73758a76e0 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_has_dropout_maxk_128.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, + true, + false, + true, + 128>(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 new file mode 100644 index 0000000000..cbefee339d --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_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, + 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 new file mode 100644 index 0000000000..84fe9cf731 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_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, + 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_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_no_dropout_maxk_128.cpp new file mode 100644 index 0000000000..7e38a00611 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_no_dropout_maxk_128.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, + true, + false, + false, + 128>(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 new file mode 100644 index 0000000000..d0f8a1131a --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_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, + 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 new file mode 100644 index 0000000000..a1eb2ed6f6 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_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, + 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_has_causal_no_local_has_bias_has_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_128.cpp new file mode 100644 index 0000000000..88a4389fa9 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_128.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, + false, + true, + true, + 128>(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 new file mode 100644 index 0000000000..58bc729e42 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_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, + 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 new file mode 100644 index 0000000000..6e8ba77772 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_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, + 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_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_128.cpp new file mode 100644 index 0000000000..f2255a6feb --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_128.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, + false, + true, + false, + 128>(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 new file mode 100644 index 0000000000..f717ee4563 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_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, + 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 new file mode 100644 index 0000000000..258b1a576e --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_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, + 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_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_128.cpp new file mode 100644 index 0000000000..c659bec52a --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_128.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, + false, + false, + true, + 128>(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 new file mode 100644 index 0000000000..d9b4631acd --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_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, + 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 new file mode 100644 index 0000000000..7796f67298 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_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, + 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_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_128.cpp new file mode 100644 index 0000000000..def077decb --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_128.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, + false, + false, + false, + 128>(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 new file mode 100644 index 0000000000..a442eda92b --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_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, + 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 new file mode 100644 index 0000000000..d79f8ffacd --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_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, + 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_no_causal_has_local_has_bias_has_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_128.cpp new file mode 100644 index 0000000000..a5be732705 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_128.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, + true, + true, + true, + 128>(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 new file mode 100644 index 0000000000..cbccdcdf0a --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_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, + 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 new file mode 100644 index 0000000000..ea329e073f --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_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, + 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_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_128.cpp new file mode 100644 index 0000000000..15c3963ef1 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_128.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, + true, + true, + false, + 128>(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 new file mode 100644 index 0000000000..614eae2111 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_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, + 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 new file mode 100644 index 0000000000..428e08cc13 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_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, + 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_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_128.cpp new file mode 100644 index 0000000000..e7c7acc1d4 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_128.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, + true, + false, + true, + 128>(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 new file mode 100644 index 0000000000..aa35336a75 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_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, + 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 new file mode 100644 index 0000000000..3070211e3f --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_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, + 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_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_128.cpp new file mode 100644 index 0000000000..9d3e1b2b18 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_128.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, + true, + false, + false, + 128>(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 new file mode 100644 index 0000000000..44abb0207f --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_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, + 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 new file mode 100644 index 0000000000..b08d8d7012 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_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, + 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_has_causal_no_local_has_bias_has_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_has_bias_has_dropout_maxk_128.cpp new file mode 100644 index 0000000000..2cd8232768 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_has_bias_has_dropout_maxk_128.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, + false, + true, + true, + 128>(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 new file mode 100644 index 0000000000..68af8bdfaf --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_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, + 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 new file mode 100644 index 0000000000..0bedd49441 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_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, + 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_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_has_bias_no_dropout_maxk_128.cpp new file mode 100644 index 0000000000..57b7ab3e74 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_has_bias_no_dropout_maxk_128.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, + false, + true, + false, + 128>(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 new file mode 100644 index 0000000000..3679db4777 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_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, + 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 new file mode 100644 index 0000000000..dd9ad9c2bf --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_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, + 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_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_no_bias_has_dropout_maxk_128.cpp new file mode 100644 index 0000000000..5d6ac84541 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_no_bias_has_dropout_maxk_128.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, + false, + false, + true, + 128>(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 new file mode 100644 index 0000000000..33d24204a9 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_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, + 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 new file mode 100644 index 0000000000..12f1013041 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_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, + 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_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_no_bias_no_dropout_maxk_128.cpp new file mode 100644 index 0000000000..cbda00a646 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_causal_no_local_no_bias_no_dropout_maxk_128.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, + false, + false, + false, + 128>(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 new file mode 100644 index 0000000000..4693f6ea05 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_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, + 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 new file mode 100644 index 0000000000..c4427d8e6e --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_has_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, + 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_no_causal_has_local_has_bias_has_dropout_maxk_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_has_bias_has_dropout_maxk_128.cpp new file mode 100644 index 0000000000..24e0a83250 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_has_bias_has_dropout_maxk_128.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, + true, + true, + true, + 128>(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 new file mode 100644 index 0000000000..00be5745df --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_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, + 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 new file mode 100644 index 0000000000..2cb1cd7e83 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_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, + 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_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_has_bias_no_dropout_maxk_128.cpp new file mode 100644 index 0000000000..05620a79a3 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_has_bias_no_dropout_maxk_128.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, + true, + true, + false, + 128>(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 new file mode 100644 index 0000000000..25b0dc246e --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_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, + 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 new file mode 100644 index 0000000000..a4b8bb18e8 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_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, + 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_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_no_bias_has_dropout_maxk_128.cpp new file mode 100644 index 0000000000..0a5e57c6af --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_no_bias_has_dropout_maxk_128.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, + true, + false, + true, + 128>(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 new file mode 100644 index 0000000000..af606140ed --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_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, + 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 new file mode 100644 index 0000000000..4b05b4e9b7 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_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, + 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_128.cpp b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_no_bias_no_dropout_maxk_128.cpp new file mode 100644 index 0000000000..c4c5c2625c --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_causal_has_local_no_bias_no_dropout_maxk_128.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, + true, + false, + false, + 128>(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 new file mode 100644 index 0000000000..b14466864c --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_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, + 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 new file mode 100644 index 0000000000..04e311ca20 --- /dev/null +++ b/example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_fp16_no_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, + false, + true, + false, + false, + 64>(HstuAttentionFwdParams& param, hipStream_t stream);