mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-06-28 18:56:59 +00:00
Remove the kHasBias==true instances to save building time
This commit is contained in:
@@ -99,7 +99,7 @@ def create_forward_instances(instance_dir: Path, headdims: List) -> None:
|
||||
(True, True),
|
||||
(False, False),
|
||||
]:
|
||||
for has_bias in [True, False]:
|
||||
for has_bias in [False]:
|
||||
for has_dropout in [False]:
|
||||
for max_k in headdims:
|
||||
fname = HSTU_FORWARD_INSTANCE_FNAME.format(
|
||||
@@ -158,7 +158,7 @@ def create_forward_instances_ref(instance_dir: Path, headdims: List) -> None:
|
||||
file.write(HSTU_COPYRIGHT_HEADER)
|
||||
file.write(forward_instance_inc)
|
||||
for max_k in headdims:
|
||||
for has_bias in [True, False]:
|
||||
for has_bias in [False]:
|
||||
for has_dropout in [False]:
|
||||
for has_causal in [True, False]:
|
||||
for use_softmax, store_lse in [
|
||||
|
||||
@@ -8,11 +8,11 @@
|
||||
|
||||
void hstu_attention_group_forward_bf16(HstuAttentionGroupFwdParams& param, hipStream_t stream)
|
||||
{
|
||||
const bool has_bias = (param.bias_ptr != nullptr);
|
||||
const bool use_causal = param.use_causal;
|
||||
bool store_lse = (param.use_softmax && param.is_training);
|
||||
|
||||
BOOL_SWITCH_3(has_bias, kHasBias, use_causal, kUseCausal, param.use_softmax, kUseSoftmax, [&] {
|
||||
constexpr bool kHasBias = false;
|
||||
BOOL_SWITCH_2(use_causal, kUseCausal, param.use_softmax, kUseSoftmax, [&] {
|
||||
HDIM_SWITCH(param.hdim_qk, param.hdim_v, MaxK, [&] {
|
||||
BOOL_SWITCH(store_lse, kStoreLSE, [&] {
|
||||
if constexpr(kUseSoftmax || !kStoreLSE)
|
||||
|
||||
@@ -8,11 +8,11 @@
|
||||
|
||||
void hstu_attention_group_forward_fp16(HstuAttentionGroupFwdParams& param, hipStream_t stream)
|
||||
{
|
||||
const bool has_bias = (param.bias_ptr != nullptr);
|
||||
const bool use_causal = param.use_causal;
|
||||
const bool store_lse = (param.use_softmax && param.is_training);
|
||||
|
||||
BOOL_SWITCH_3(has_bias, kHasBias, use_causal, kUseCausal, param.use_softmax, kUseSoftmax, [&] {
|
||||
constexpr bool kHasBias = false;
|
||||
BOOL_SWITCH_2(use_causal, kUseCausal, param.use_softmax, kUseSoftmax, [&] {
|
||||
HDIM_SWITCH(param.hdim_qk, param.hdim_v, MaxK, [&] {
|
||||
BOOL_SWITCH(store_lse, kStoreLSE, [&] {
|
||||
if constexpr(kUseSoftmax || !kStoreLSE)
|
||||
|
||||
@@ -9,11 +9,11 @@
|
||||
|
||||
void hstu_attention_no_group_forward_bf16(HstuAttentionNoGroupFwdParams& param, hipStream_t stream)
|
||||
{
|
||||
const bool has_bias = (param.bias_ptr != nullptr);
|
||||
const bool use_causal = param.use_causal;
|
||||
bool store_lse = (param.use_softmax && param.is_training);
|
||||
|
||||
BOOL_SWITCH_3(has_bias, kHasBias, use_causal, kUseCausal, param.use_softmax, kUseSoftmax, [&] {
|
||||
constexpr bool kHasBias = false;
|
||||
BOOL_SWITCH_2(use_causal, kUseCausal, param.use_softmax, kUseSoftmax, [&] {
|
||||
HDIM_SWITCH(param.hdim_qk, param.hdim_v, MaxK, [&] {
|
||||
BOOL_SWITCH(store_lse, kStoreLSE, [&] {
|
||||
if constexpr(kUseSoftmax || !kStoreLSE)
|
||||
|
||||
@@ -9,11 +9,11 @@
|
||||
|
||||
void hstu_attention_no_group_forward_fp16(HstuAttentionNoGroupFwdParams& param, hipStream_t stream)
|
||||
{
|
||||
const bool has_bias = (param.bias_ptr != nullptr);
|
||||
const bool use_causal = param.use_causal;
|
||||
bool store_lse = (param.use_softmax && param.is_training);
|
||||
|
||||
BOOL_SWITCH_3(has_bias, kHasBias, use_causal, kUseCausal, param.use_softmax, kUseSoftmax, [&] {
|
||||
constexpr bool kHasBias = false;
|
||||
BOOL_SWITCH_2(use_causal, kUseCausal, param.use_softmax, kUseSoftmax, [&] {
|
||||
HDIM_SWITCH(param.hdim_qk, param.hdim_v, MaxK, [&] {
|
||||
BOOL_SWITCH(store_lse, kStoreLSE, [&] {
|
||||
if constexpr(kUseSoftmax || !kStoreLSE)
|
||||
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -10,60 +10,6 @@
|
||||
#include "hstu_attention_batched_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
@@ -118,60 +64,6 @@ extern template void run_batched_forward_dispatch<
|
||||
false,
|
||||
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
@@ -226,60 +118,6 @@ extern template void run_batched_forward_dispatch<
|
||||
false,
|
||||
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
@@ -334,60 +172,6 @@ extern template void run_batched_forward_dispatch<
|
||||
false,
|
||||
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -10,60 +10,6 @@
|
||||
#include "hstu_attention_batched_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
@@ -118,60 +64,6 @@ extern template void run_batched_forward_dispatch<
|
||||
false,
|
||||
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
@@ -226,60 +118,6 @@ extern template void run_batched_forward_dispatch<
|
||||
false,
|
||||
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
@@ -334,60 +172,6 @@ extern template void run_batched_forward_dispatch<
|
||||
false,
|
||||
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_batched_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionNoGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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_group_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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_group_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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_group_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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_group_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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_group_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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_group_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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_group_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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_group_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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_group_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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_group_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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_group_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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_group_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -10,60 +10,6 @@
|
||||
#include "hstu_attention_group_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
@@ -118,60 +64,6 @@ extern template void run_group_forward_dispatch<
|
||||
false,
|
||||
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
@@ -226,60 +118,6 @@ extern template void run_group_forward_dispatch<
|
||||
false,
|
||||
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
@@ -334,60 +172,6 @@ extern template void run_group_forward_dispatch<
|
||||
false,
|
||||
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
true,
|
||||
|
||||
@@ -1,20 +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_group_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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_group_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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_group_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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_group_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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_group_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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_group_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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_group_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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_group_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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_group_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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_group_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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_group_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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_group_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::bf16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -10,60 +10,6 @@
|
||||
#include "hstu_attention_group_forward_dispatch.hpp"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
@@ -118,60 +64,6 @@ extern template void run_group_forward_dispatch<
|
||||
false,
|
||||
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
@@ -226,60 +118,6 @@ extern template void run_group_forward_dispatch<
|
||||
false,
|
||||
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
@@ -334,60 +172,6 @@ extern template void run_group_forward_dispatch<
|
||||
false,
|
||||
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
true,
|
||||
true,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
|
||||
extern template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
true,
|
||||
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
false,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
96>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
128>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
256>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
@@ -1,20 +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"
|
||||
#include "hstu_attention_params.hpp"
|
||||
|
||||
template void run_group_forward_dispatch<
|
||||
ck_tile::fp16_t,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
true,
|
||||
false,
|
||||
64>(HstuAttentionGroupFwdParams& param, hipStream_t stream);
|
||||
Some files were not shown because too many files have changed in this diff Show More
Reference in New Issue
Block a user