ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-15 10:37:44 +00:00
Code Issues Packages Projects Releases Wiki Activity

Revert "Disable support of hdim64 amnd hdim256 for quick compiling and testing"

Browse Source 
This reverts commit ed062f93e2.
This commit is contained in:
Qianfeng Zhang
2025-07-17 09:00:41 +00:00
parent ed062f93e2
commit fed1474e4f
134 changed files with 3340 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
example/ck_tile/18_hstu_attention/generate_instances.py
View File

@@ -161,7 +161,7 @@ def create_forward_instances_ref(instance_dir: Path, headdims: List) -> None:
file.write(forward_instance)
if __name__ == "__main__":
headdims_fwd = [128]
headdims_fwd = [64, 128, 256]
this_dir = os.path.dirname(__file__)
output_dir = Path(this_dir) / "instances"

12
example/ck_tile/18_hstu_attention/hstu_attention_hdim_switch.hpp
View File

@@ -8,11 +8,21 @@
#define HDIM_SWITCH(HDIM_1, HDIM_2, CONST_NAME, ...) \
[&] { \
if(HDIM_1 <= 128 && HDIM_2 <= 128) \
if(HDIM_1 <= 64 && HDIM_2 <= 64) \
{ \
constexpr ck_tile::index_t CONST_NAME = 64; \
__VA_ARGS__(); \
} \
else if(HDIM_1 <= 128 && HDIM_2 <= 128) \
{ \
constexpr ck_tile::index_t CONST_NAME = 128; \
__VA_ARGS__(); \
} \
else if(HDIM_1 <= 256 && HDIM_2 <= 256) \
{ \
constexpr ck_tile::index_t CONST_NAME = 256; \
__VA_ARGS__(); \
} \
else \
{ \
throw std::runtime_error("Head-dim sizes not supported!"); \

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
true,
256>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
true,
64>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
false,
256>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
false,
64>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
true,
256>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
true,
64>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
false,
256>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
false,
64>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

128
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_instances_ref.hpp
View File

@@ -13,6 +13,70 @@
#include <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
true,
64>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
true,
64>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
false,
64>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
false,
64>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
true,
64>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
true,
64>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
false,
64>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
false,
64>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
@@ -76,3 +140,67 @@ extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
false,
false,
128>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
true,
256>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
true,
256>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
false,
256>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
false,
256>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
true,
256>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
true,
256>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
false,
256>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
false,
256>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
true,
256>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
true,
64>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
false,
256>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
false,
64>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
true,
256>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
true,
64>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
false,
256>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
false,
64>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
true,
true,
256>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
true,
true,
64>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
true,
false,
256>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
true,
false,
64>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
false,
true,
256>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
false,
true,
64>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
false,
false,
256>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
false,
false,
64>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_has_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_has_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_no_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_has_bias_no_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_has_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_has_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_no_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_no_local_no_bias_no_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#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);

128
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_instances_ref.hpp
View File

@@ -13,6 +13,70 @@
#include <ck_tile/core/numeric/bfloat16.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
true,
true,
64>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
true,
true,
64>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
true,
false,
64>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
true,
false,
64>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
false,
true,
64>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
false,
true,
64>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
false,
false,
64>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
false,
false,
64>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
@@ -76,3 +140,67 @@ extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
false,
false,
128>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
true,
true,
256>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
true,
true,
256>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
true,
false,
256>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
true,
false,
256>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
false,
true,
256>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
false,
true,
256>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
false,
false,
256>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
false,
false,
256>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_has_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_has_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_no_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_has_bias_no_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_has_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_has_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_no_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_has_local_no_bias_no_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
true,
true,
256>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
true,
true,
64>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
true,
false,
256>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
true,
false,
64>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
false,
true,
256>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
false,
true,
64>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
false,
false,
256>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_no_local_no_bias_no_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/bfloat16.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
template void run_batched_forward_causal_local_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
false,
false,
64>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_jagged_forward_dispatch.hpp"
template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
true,
256>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_has_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_jagged_forward_dispatch.hpp"
template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
true,
64>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_jagged_forward_dispatch.hpp"
template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
false,
256>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_has_bias_no_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_jagged_forward_dispatch.hpp"
template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
false,
64>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_jagged_forward_dispatch.hpp"
template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
true,
256>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_has_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_jagged_forward_dispatch.hpp"
template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
true,
64>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_jagged_forward_dispatch.hpp"
template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
false,
256>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_has_local_no_bias_no_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_jagged_forward_dispatch.hpp"
template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
false,
64>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_has_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_has_bias_no_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_has_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_has_causal_no_local_no_bias_no_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

128
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_instances_ref.hpp
View File

@@ -13,6 +13,70 @@
#include <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_jagged_forward_dispatch.hpp"
extern template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
true,
64>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
true,
64>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
false,
64>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
false,
64>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
true,
64>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
true,
64>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
false,
64>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
false,
64>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
@@ -76,3 +140,67 @@ extern template void run_jagged_forward_causal_local_bias_dropout_dispatch<
false,
false,
128>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
true,
256>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
true,
256>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
false,
256>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
false,
256>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
true,
256>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
true,
256>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
false,
256>(HstuAttentionFwdParams& param, hipStream_t stream);
extern template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
false,
256>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_has_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_has_bias_no_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_has_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_has_local_no_bias_no_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#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);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_jagged_forward_dispatch.hpp"
template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
true,
256>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_has_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_jagged_forward_dispatch.hpp"
template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
true,
64>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_jagged_forward_dispatch.hpp"
template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
false,
256>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_has_bias_no_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_jagged_forward_dispatch.hpp"
template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
false,
64>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_jagged_forward_dispatch.hpp"
template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
true,
256>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_has_dropout_maxk_64.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_jagged_forward_dispatch.hpp"
template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
true,
64>(HstuAttentionFwdParams& param, hipStream_t stream);

22
example/ck_tile/18_hstu_attention/instances/hstu_attention_jagged_forward_bf16_no_causal_no_local_no_bias_no_dropout_maxk_256.cpp Normal file
View File

@@ -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 <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_jagged_forward_dispatch.hpp"
template void run_jagged_forward_causal_local_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
false,
256>(HstuAttentionFwdParams& param, hipStream_t stream);

Some files were not shown because too many files have changed in this diff Show More