From a812b044ae4f91abf0e0028592d248b5d4d87fa3 Mon Sep 17 00:00:00 2001 From: Linjun-AMD Date: Mon, 2 Feb 2026 01:34:23 -0600 Subject: [PATCH] updated fmha_args Signed-off-by: Linjun-AMD --- example/ck_tile/01_fmha/fmha_fwd.hpp | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) diff --git a/example/ck_tile/01_fmha/fmha_fwd.hpp b/example/ck_tile/01_fmha/fmha_fwd.hpp index aedbb0e17c..65c280d42d 100644 --- a/example/ck_tile/01_fmha/fmha_fwd.hpp +++ b/example/ck_tile/01_fmha/fmha_fwd.hpp @@ -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;