ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-06-06 15:54:31 +00:00
Code Issues Packages Projects Releases Wiki Activity

Fix compilation errors

Browse Source
This commit is contained in:
PoYen, Chen
2024-08-08 08:26:55 +00:00
parent 247e135cfc
commit 9d9c5a6c24
3 changed files with 7 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
example/ck_tile/01_fmha/codegen/ops/fmha_fwd_splitkv.py
View File

@@ -220,7 +220,7 @@ float fmha_fwd_splitkv_(const ck_tile::stream_config& s, fmha_fwd_splitkv_args a
);
}}
float fmha_fwd_splitkv(fmha_fwd_traits t, fmha_fwd_splitkv_args a, const ck_tile::stream_config& s){{
float fmha_fwd_splitkv(fmha_fwd_splitkv_traits t, fmha_fwd_splitkv_args a, const ck_tile::stream_config& s){{
float r = -1;
{F_dispatch}
return r;

6
example/ck_tile/01_fmha/fmha_fwd.cpp
View File

@@ -989,12 +989,12 @@ bool run(const ck_tile::ArgParser& arg_parser)
const float fwd_ave_time = [&] {
#if CK_TILE_FMHA_FWD_SPLITKV_API
if(1 < args.num_splits || args.block_table_ptr != nullptr)
if(1 < num_splits || 0 < page_block_size)
{
return fmha_fwd_splitkv(traits, args, config);
return fmha_fwd_splitkv(fmha_traits, fmha_args, stream_config);
}
#endif
return fmha_fwd(traits, args, config);
return fmha_fwd(fmha_traits, fmha_args, stream_config);
}();
if(appendkv_ave_time < 0 || fwd_ave_time < 0)

4
example/ck_tile/01_fmha/fmha_fwd.hpp
View File

@@ -745,7 +745,9 @@ struct fmha_fwd_traits
// TODO: padding check is inside this api
};
float fmha_fwd(fmha_fwd_traits, fmha_fwd_args, const ck_tile::stream_config&);
float fmha_fwd_splitkv(fmha_fwd_traits, fmha_fwd_args, const ck_tile::stream_config&);
using fmha_fwd_splitkv_traits = fmha_fwd_traits;
float fmha_fwd_splitkv(fmha_fwd_splitkv_traits, fmha_fwd_splitkv_args, const ck_tile::stream_config&);
struct fmha_fwd_appendkv_traits
{