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

Remove dropout=true instances to reduce compiling-time

Browse Source
This commit is contained in:
Qianfeng Zhang
2026-04-07 09:38:18 +00:00
parent 348c3e05be
commit a95f64601d
203 changed files with 42 additions and 5032 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -82,19 +82,22 @@ MODE_GROUP_OR_NOT_MAP = {
"group": "Group",
}
def create_forward_instances(instance_dir: Path, headdims: List) -> None:
for mode in ["batched", "jagged", "group"]:
for dtype in ["fp16", "bf16"]:
for has_causal in [True, False]:
for use_softmax in [True, False]:
for has_bias in [True, False]:
for has_dropout in [True, False]:
for has_dropout in [False]:
for max_k in headdims:
fname = HSTU_FORWARD_INSTANCE_FNAME.format(
mode=mode,
dtype_str=dtype,
has_or_no_causal_str=BOOL_MAP_CAUSAL[has_causal],
use_softmax_or_not_str=BOOL_MAP_SOFTMAX[use_softmax],
use_softmax_or_not_str=BOOL_MAP_SOFTMAX[
use_softmax
],
has_or_no_bias_str=BOOL_MAP_BIAS[has_bias],
has_or_no_dropout_str=BOOL_MAP_DROPOUT[has_dropout],
max_k_str=INT_MAP_MAX_K[max_k],
@@ -105,22 +108,24 @@ def create_forward_instances(instance_dir: Path, headdims: List) -> None:
dtype_file=TYPE_FNAME_MAP[dtype],
)
)
forward_instance = HSTU_FORWARD_INSTANCE_TEMPLATE.format(
extern="",
mode=mode,
dtype=TYPE_CTYPE_MAP[dtype],
has_causal=BOOL_MAP[has_causal],
use_softmax=BOOL_MAP[use_softmax],
has_bias=BOOL_MAP[has_bias],
has_dropout=BOOL_MAP[has_dropout],
max_k=max_k,
group_or_not=MODE_GROUP_OR_NOT_MAP[mode],
forward_instance = (
HSTU_FORWARD_INSTANCE_TEMPLATE.format(
extern="",
mode=mode,
dtype=TYPE_CTYPE_MAP[dtype],
has_causal=BOOL_MAP[has_causal],
use_softmax=BOOL_MAP[use_softmax],
has_bias=BOOL_MAP[has_bias],
has_dropout=BOOL_MAP[has_dropout],
max_k=max_k,
group_or_not=MODE_GROUP_OR_NOT_MAP[mode],
)
)
(instance_dir / fname).write_text(
HSTU_COPYRIGHT_HEADER
+ forward_instance_inc
+ forward_instance
)
HSTU_COPYRIGHT_HEADER
+ forward_instance_inc
+ forward_instance
)
def create_forward_instances_ref(instance_dir: Path, headdims: List) -> None:
@@ -141,7 +146,7 @@ def create_forward_instances_ref(instance_dir: Path, headdims: List) -> None:
file.write(forward_instance_inc)
for max_k in headdims:
for has_bias in [True, False]:
for has_dropout in [True, False]:
for has_dropout in [False]:
for has_causal in [True, False]:
for use_softmax in [True, False]:
forward_instance = (
@@ -156,9 +161,10 @@ def create_forward_instances_ref(instance_dir: Path, headdims: List) -> None:
max_k=max_k,
group_or_not=MODE_GROUP_OR_NOT_MAP[mode],
)
)
)
file.write(forward_instance)
if __name__ == "__main__":
headdims_fwd = [64, 96, 128, 256]

9
example/ck_tile/18_hstu_attention/hstu_attention_group_forward_bf16.cpp
View File

@@ -12,17 +12,16 @@
void hstu_attention_group_forward_bf16(HstuAttentionGroupFwdParams& param, hipStream_t stream)
{
const bool has_dropout = (param.p_drop > 0.0f);
const bool has_bias = (param.bias_ptr != nullptr);
const bool use_causal = param.use_causal;
BOOL_SWITCH_3(has_bias, kHasBias, has_dropout, kHasDropout, use_causal, kUseCausal, [&] {
const bool has_bias = (param.bias_ptr != nullptr);
const bool use_causal = param.use_causal;
BOOL_SWITCH_2(has_bias, kHasBias, use_causal, kUseCausal, [&] {
HDIM_SWITCH(param.hdim_qk, param.hdim_v, MaxK, [&] {
BOOL_SWITCH(param.use_softmax, kUseSoftmax, [&] {
run_group_forward_causal_softmax_bias_dropout_dispatch<ck_tile::bf16_t,
kUseCausal,
kUseSoftmax,
kHasBias,
kHasDropout,
false, // kHasDropout
MaxK>(param, stream);
});
});

9
example/ck_tile/18_hstu_attention/hstu_attention_group_forward_fp16.cpp
View File

@@ -12,18 +12,17 @@
void hstu_attention_group_forward_fp16(HstuAttentionGroupFwdParams& param, hipStream_t stream)
{
const bool has_dropout = (param.p_drop > 0.0f);
const bool has_bias = (param.bias_ptr != nullptr);
const bool use_causal = param.use_causal;
const bool has_bias = (param.bias_ptr != nullptr);
const bool use_causal = param.use_causal;
BOOL_SWITCH_3(has_bias, kHasBias, has_dropout, kHasDropout, use_causal, kUseCausal, [&] {
BOOL_SWITCH_2(has_bias, kHasBias, use_causal, kUseCausal, [&] {
HDIM_SWITCH(param.hdim_qk, param.hdim_v, MaxK, [&] {
BOOL_SWITCH(param.use_softmax, kUseSoftmax, [&] {
run_group_forward_causal_softmax_bias_dropout_dispatch<ck_tile::fp16_t,
kUseCausal,
kUseSoftmax,
kHasBias,
kHasDropout,
false, // kHasDropout
MaxK>(param, stream);
});
});

11
example/ck_tile/18_hstu_attention/hstu_attention_no_group_forward_bf16.cpp
View File

@@ -14,10 +14,9 @@
void hstu_attention_no_group_forward_bf16(HstuAttentionNoGroupFwdParams& param, hipStream_t stream)
{
const bool has_dropout = (param.p_drop > 0.0f);
const bool has_bias = (param.bias_ptr != nullptr);
const bool use_causal = param.use_causal;
BOOL_SWITCH_3(has_bias, kHasBias, has_dropout, kHasDropout, use_causal, kUseCausal, [&] {
const bool has_bias = (param.bias_ptr != nullptr);
const bool use_causal = param.use_causal;
BOOL_SWITCH_2(has_bias, kHasBias, use_causal, kUseCausal, [&] {
HDIM_SWITCH(param.hdim_qk, param.hdim_v, MaxK, [&] {
BOOL_SWITCH(param.use_softmax, kUseSoftmax, [&] {
if(param.is_jagged)
@@ -25,14 +24,14 @@ void hstu_attention_no_group_forward_bf16(HstuAttentionNoGroupFwdParams& param,
kUseCausal,
kUseSoftmax,
kHasBias,
kHasDropout,
false, // kHasDropout
MaxK>(param, stream);
else
run_batched_forward_causal_softmax_bias_dropout_dispatch<ck_tile::bf16_t,
kUseCausal,
kUseSoftmax,
kHasBias,
kHasDropout,
false, // kHasDropout
MaxK>(param, stream);
});
});

11
example/ck_tile/18_hstu_attention/hstu_attention_no_group_forward_fp16.cpp
View File

@@ -14,10 +14,9 @@
void hstu_attention_no_group_forward_fp16(HstuAttentionNoGroupFwdParams& param, hipStream_t stream)
{
const bool has_dropout = (param.p_drop > 0.0f);
const bool has_bias = (param.bias_ptr != nullptr);
const bool use_causal = param.use_causal;
BOOL_SWITCH_3(has_bias, kHasBias, has_dropout, kHasDropout, use_causal, kUseCausal, [&] {
const bool has_bias = (param.bias_ptr != nullptr);
const bool use_causal = param.use_causal;
BOOL_SWITCH_2(has_bias, kHasBias, use_causal, kUseCausal, [&] {
HDIM_SWITCH(param.hdim_qk, param.hdim_v, MaxK, [&] {
BOOL_SWITCH(param.use_softmax, kUseSoftmax, [&] {
if(param.is_jagged)
@@ -25,14 +24,14 @@ void hstu_attention_no_group_forward_fp16(HstuAttentionNoGroupFwdParams& param,
kUseCausal,
kUseSoftmax,
kHasBias,
kHasDropout,
false, // kHasDropout
MaxK>(param, stream);
else
run_batched_forward_causal_softmax_bias_dropout_dispatch<ck_tile::fp16_t,
kUseCausal,
kUseSoftmax,
kHasBias,
kHasDropout,
false, // kHasDropout
MaxK>(param, stream);
});
});

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_has_bias_has_dropout_maxk_128.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
true,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_has_bias_has_dropout_maxk_256.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
true,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_has_bias_has_dropout_maxk_64.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
true,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_has_bias_has_dropout_maxk_96.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
true,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_no_bias_has_dropout_maxk_128.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
false,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_no_bias_has_dropout_maxk_256.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
false,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_no_bias_has_dropout_maxk_64.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
false,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_false_no_bias_has_dropout_maxk_96.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
false,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_has_bias_has_dropout_maxk_128.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_has_bias_has_dropout_maxk_256.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_has_bias_has_dropout_maxk_64.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_has_bias_has_dropout_maxk_96.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_no_bias_has_dropout_maxk_128.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_no_bias_has_dropout_maxk_256.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_no_bias_has_dropout_maxk_64.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_has_causal_softmax_true_no_bias_has_dropout_maxk_96.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

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

@@ -9,38 +9,6 @@
#include <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
true,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
true,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
@@ -73,38 +41,6 @@ extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
false,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
false,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
false,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
@@ -137,38 +73,6 @@ extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
false,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
true,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
true,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
@@ -201,38 +105,6 @@ extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
false,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
false,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
false,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
@@ -265,38 +137,6 @@ extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
false,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
true,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
true,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
@@ -329,38 +169,6 @@ extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
false,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
false,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
false,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
@@ -393,38 +201,6 @@ extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
false,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
true,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
true,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
@@ -457,38 +233,6 @@ extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
false,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
false,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
false,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_has_bias_has_dropout_maxk_128.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_has_bias_has_dropout_maxk_256.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_has_bias_has_dropout_maxk_64.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_has_bias_has_dropout_maxk_96.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_no_bias_has_dropout_maxk_128.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_no_bias_has_dropout_maxk_256.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_no_bias_has_dropout_maxk_64.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_false_no_bias_has_dropout_maxk_96.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_has_bias_has_dropout_maxk_128.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
true,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_has_bias_has_dropout_maxk_256.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
true,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_has_bias_has_dropout_maxk_64.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
true,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_has_bias_has_dropout_maxk_96.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
true,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_no_bias_has_dropout_maxk_128.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
false,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_no_bias_has_dropout_maxk_256.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
false,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_no_bias_has_dropout_maxk_64.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
false,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_bf16_no_causal_softmax_true_no_bias_has_dropout_maxk_96.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
false,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_has_bias_has_dropout_maxk_128.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
false,
true,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_has_bias_has_dropout_maxk_256.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
false,
true,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_has_bias_has_dropout_maxk_64.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
false,
true,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_has_bias_has_dropout_maxk_96.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
false,
true,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_no_bias_has_dropout_maxk_128.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
false,
false,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_no_bias_has_dropout_maxk_256.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
false,
false,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_no_bias_has_dropout_maxk_64.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
false,
false,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_false_no_bias_has_dropout_maxk_96.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
false,
false,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_has_bias_has_dropout_maxk_128.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
true,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_has_bias_has_dropout_maxk_256.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
true,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_has_bias_has_dropout_maxk_64.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
true,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_has_bias_has_dropout_maxk_96.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
true,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_no_bias_has_dropout_maxk_128.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
false,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_no_bias_has_dropout_maxk_256.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
false,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_no_bias_has_dropout_maxk_64.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
false,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_has_causal_softmax_true_no_bias_has_dropout_maxk_96.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
false,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

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

@@ -9,38 +9,6 @@
#include <ck_tile/core/numeric/bfloat16.hpp>
#include "hstu_attention_batched_forward_dispatch.hpp"
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
true,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
false,
true,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
true,
true,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
true,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
@@ -73,38 +41,6 @@ extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
false,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
false,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
false,
false,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
true,
false,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
false,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
@@ -137,38 +73,6 @@ extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
false,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
true,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
false,
true,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
true,
true,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
true,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
@@ -201,38 +105,6 @@ extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
false,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
false,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
false,
false,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
true,
false,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
false,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
@@ -265,38 +137,6 @@ extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
false,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
true,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
false,
true,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
true,
true,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
true,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
@@ -329,38 +169,6 @@ extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
false,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
false,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
false,
false,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
true,
false,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
false,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
@@ -393,38 +201,6 @@ extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
false,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
true,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
false,
true,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
true,
true,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
true,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
@@ -457,38 +233,6 @@ extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
false,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
true,
false,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,
false,
false,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
true,
false,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
false,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
extern template void run_batched_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
true,

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_has_bias_has_dropout_maxk_128.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
true,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_has_bias_has_dropout_maxk_256.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
true,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_has_bias_has_dropout_maxk_64.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
true,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_has_bias_has_dropout_maxk_96.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
true,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_no_bias_has_dropout_maxk_128.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
false,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_no_bias_has_dropout_maxk_256.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
false,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_no_bias_has_dropout_maxk_64.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
false,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_false_no_bias_has_dropout_maxk_96.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
false,
false,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_has_bias_has_dropout_maxk_128.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
true,
true,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_has_bias_has_dropout_maxk_256.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
true,
true,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_has_bias_has_dropout_maxk_64.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
true,
true,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_has_bias_has_dropout_maxk_96.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
true,
true,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_no_bias_has_dropout_maxk_128.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
true,
false,
true,
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_no_bias_has_dropout_maxk_256.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
true,
false,
true,
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_no_bias_has_dropout_maxk_64.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
true,
false,
true,
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_batched_forward_fp16_no_causal_softmax_true_no_bias_has_dropout_maxk_96.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_softmax_bias_dropout_dispatch<
ck_tile::fp16_t,
false,
true,
false,
true,
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_false_has_bias_has_dropout_maxk_128.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
true,
true,
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_false_has_bias_has_dropout_maxk_256.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
true,
true,
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_false_has_bias_has_dropout_maxk_64.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
true,
true,
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_false_has_bias_has_dropout_maxk_96.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
true,
true,
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_false_no_bias_has_dropout_maxk_128.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
false,
true,
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_false_no_bias_has_dropout_maxk_256.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
false,
true,
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_false_no_bias_has_dropout_maxk_64.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
false,
true,
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_false_no_bias_has_dropout_maxk_96.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
false,
true,
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_has_bias_has_dropout_maxk_128.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
true,
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_has_bias_has_dropout_maxk_256.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
true,
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_has_bias_has_dropout_maxk_64.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
true,
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_has_bias_has_dropout_maxk_96.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
true,
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_no_bias_has_dropout_maxk_128.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
true,
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_no_bias_has_dropout_maxk_256.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
true,
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_no_bias_has_dropout_maxk_64.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
true,
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_has_causal_softmax_true_no_bias_has_dropout_maxk_96.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
true,
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

256
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_instances_ref.hpp
View File

@@ -9,38 +9,6 @@
#include <ck_tile/core/numeric/half.hpp>
#include "hstu_attention_group_forward_dispatch.hpp"
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
true,
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
true,
true,
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
true,
true,
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
true,
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
@@ -73,38 +41,6 @@ extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
false,
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
true,
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
false,
true,
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
false,
true,
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
true,
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
@@ -137,38 +73,6 @@ extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
false,
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
true,
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
true,
true,
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
true,
true,
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
true,
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
@@ -201,38 +105,6 @@ extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
false,
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
true,
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
false,
true,
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
false,
true,
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
true,
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
@@ -265,38 +137,6 @@ extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
false,
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
true,
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
true,
true,
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
true,
true,
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
true,
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
@@ -329,38 +169,6 @@ extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
false,
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
true,
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
false,
true,
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
false,
true,
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
true,
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
@@ -393,38 +201,6 @@ extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
false,
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
true,
true,
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
true,
true,
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
true,
true,
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
true,
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
@@ -457,38 +233,6 @@ extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
false,
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
true,
false,
true,
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,
false,
false,
true,
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
false,
true,
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
true,
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
extern template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
true,

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_false_has_bias_has_dropout_maxk_128.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
true,
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_false_has_bias_has_dropout_maxk_256.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
true,
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_false_has_bias_has_dropout_maxk_64.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
true,
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_false_has_bias_has_dropout_maxk_96.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
true,
true,
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_false_no_bias_has_dropout_maxk_128.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
true,
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_false_no_bias_has_dropout_maxk_256.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
true,
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_false_no_bias_has_dropout_maxk_64.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
true,
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_false_no_bias_has_dropout_maxk_96.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
false,
false,
true,
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_has_bias_has_dropout_maxk_128.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
true,
true,
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_has_bias_has_dropout_maxk_256.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
true,
true,
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_has_bias_has_dropout_maxk_64.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
true,
true,
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

18
example/ck_tile/18_hstu_attention/instances/hstu_attention_group_forward_bf16_no_causal_softmax_true_has_bias_has_dropout_maxk_96.cpp
View File

@@ -1,18 +0,0 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
// 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_group_forward_dispatch.hpp"
template void run_group_forward_causal_softmax_bias_dropout_dispatch<
ck_tile::bf16_t,
false,
true,
true,
true,
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);

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