ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-13 18:51:13 +00:00
Code Issues Packages Projects Releases Wiki Activity

updated fmha_args

Browse Source 
Signed-off-by: Linjun-AMD <Jun.Lin@amd.com>
This commit is contained in:
Linjun-AMD
2026-02-02 01:34:23 -06:00
parent 93ef0b4fad
commit a812b044ae
1 changed files with 4 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -232,7 +232,7 @@ struct fmha_fwd_args
// array [batch + 1]. (Used with padding)
const void* block_scale_seqstart_q_ptr;
const void* block_scale_seqstart_k_ptr;
const void* sink_ptr;
const void* sink_ptr = nullptr;
ck_tile::index_t seqlen_q;
ck_tile::index_t seqlen_k;
@@ -329,7 +329,7 @@ struct fmha_fwd_pagedkv_args
const void* seqstart_q_ptr;
const void* seqstart_k_ptr;
const void* seqlen_k_ptr;
const void* sink_ptr;
const void* sink_ptr = nullptr;
ck_tile::index_t seqlen_q;
ck_tile::index_t seqlen_k;
@@ -413,7 +413,7 @@ struct fmha_fwd_splitkv_args
const void* seqstart_q_ptr;
const void* seqstart_k_ptr;
const void* seqlen_k_ptr;
const void* sink_ptr;
const void* sink_ptr = nullptr;
ck_tile::index_t seqlen_q;
ck_tile::index_t seqlen_k;
@@ -490,7 +490,6 @@ struct fmha_fwd_appendkv_args
ck_tile::index_t page_block_size; // only used if 'block_table_ptr' is not nullptr
const void* cache_batch_idx; // only used if block_table_ptr is nullptr -> batch mode (kvcache)
const void* sink_ptr;
ck_tile::index_t stride_q;
ck_tile::index_t stride_k;
@@ -534,7 +533,7 @@ struct fmha_batch_prefill_args
// 1) +
// kargs.kv_last_page_lens[b]
const void* seqstart_q_ptr;
const void* sink_ptr;
const void* sink_ptr = nullptr;
ck_tile::index_t seqlen_q;
ck_tile::index_t seqlen_k;