Skip to content

GitLab

Commit 289d5eb0 authored by Po Yen Chen's avatar Po Yen Chen
Browse files

Re-format headers

parent 612a35d6
Hide whitespace changes
Inline Side-by-side
Showing with 84 additions and 82 deletions
example/ck_tile/01_fmha/fmha_fwd.hpp
View file @ 289d5eb0
...@@ -412,91 +412,93 @@ auto fmha_fwd_splitkv_create_kargs_and_grids(fmha_fwd_splitkv_args args) ...@@ -412,91 +412,93 @@ auto fmha_fwd_splitkv_create_kargs_and_grids(fmha_fwd_splitkv_args args)
// create group mode kernel arguments // create group mode kernel arguments
if constexpr(Kernel::kIsGroupMode) if constexpr(Kernel::kIsGroupMode)
{ {
return Kernel::MakeKargs(args.q_ptr, return Kernel::MakeKargs(
args.k_ptr, args.q_ptr,
args.v_ptr, args.k_ptr,
args.bias_ptr, args.v_ptr,
(1 < args.num_splits ? args.lse_acc_ptr : args.lse_ptr), args.bias_ptr,
(1 < args.num_splits ? args.o_acc_ptr : args.o_ptr), (1 < args.num_splits ? args.lse_acc_ptr : args.lse_ptr),
args.batch, (1 < args.num_splits ? args.o_acc_ptr : args.o_ptr),
args.seqstart_q_ptr, args.batch,
args.seqstart_k_ptr, args.seqstart_q_ptr,
args.seqlen_k_ptr, args.seqstart_k_ptr,
args.hdim_q, args.seqlen_k_ptr,
args.hdim_v, args.hdim_q,
args.nhead_q, args.hdim_v,
args.nhead_q / args.nhead_k, args.nhead_q,
args.num_splits, args.nhead_q / args.nhead_k,
args.block_table_ptr, args.num_splits,
args.batch_stride_block_table, args.block_table_ptr,
args.page_block_size, args.batch_stride_block_table,
args.is_gappy, args.page_block_size,
args.scale_s, args.is_gappy,
args.scale_p, args.scale_s,
args.stride_q, args.scale_p,
args.stride_k, args.stride_q,
args.stride_v, args.stride_k,
args.stride_bias, args.stride_v,
(1 < args.num_splits ? args.stride_o_acc : args.stride_o), args.stride_bias,
args.nhead_stride_q, (1 < args.num_splits ? args.stride_o_acc : args.stride_o),
args.nhead_stride_k, args.nhead_stride_q,
args.nhead_stride_v, args.nhead_stride_k,
args.nhead_stride_bias, args.nhead_stride_v,
(1 < args.num_splits ? args.nhead_stride_lse_acc : args.nhead_stride_lse), args.nhead_stride_bias,
(1 < args.num_splits ? args.nhead_stride_o_acc : args.nhead_stride_o), (1 < args.num_splits ? args.nhead_stride_lse_acc : args.nhead_stride_lse),
args.batch_stride_k, // only used for paged-kvcache (1 < args.num_splits ? args.nhead_stride_o_acc : args.nhead_stride_o),
args.batch_stride_v, // only used for paged-kvcache args.batch_stride_k, // only used for paged-kvcache
(1 < args.num_splits ? args.split_stride_lse_acc : 0), args.batch_stride_v, // only used for paged-kvcache
(1 < args.num_splits ? args.split_stride_o_acc : 0), (1 < args.num_splits ? args.split_stride_lse_acc : 0),
args.window_size_left, (1 < args.num_splits ? args.split_stride_o_acc : 0),
args.window_size_right, args.window_size_left,
args.mask_type); args.window_size_right,
args.mask_type);
} }
else else
{ // create batch mode kernel arguments { // create batch mode kernel arguments
return Kernel::MakeKargs(args.q_ptr, return Kernel::MakeKargs(
args.k_ptr, args.q_ptr,
args.v_ptr, args.k_ptr,
args.bias_ptr, args.v_ptr,
(1 < args.num_splits ? args.lse_acc_ptr : args.lse_ptr), args.bias_ptr,
(1 < args.num_splits ? args.o_acc_ptr : args.o_ptr), (1 < args.num_splits ? args.lse_acc_ptr : args.lse_ptr),
args.batch, (1 < args.num_splits ? args.o_acc_ptr : args.o_ptr),
args.seqlen_q, args.batch,
args.seqlen_k, args.seqlen_q,
args.seqlen_k_ptr, args.seqlen_k,
args.hdim_q, args.seqlen_k_ptr,
args.hdim_v, args.hdim_q,
args.nhead_q, args.hdim_v,
args.nhead_q / args.nhead_k, args.nhead_q,
args.num_splits, args.nhead_q / args.nhead_k,
args.block_table_ptr, args.num_splits,
args.batch_stride_block_table, args.block_table_ptr,
args.page_block_size, args.batch_stride_block_table,
args.cache_batch_idx, args.page_block_size,
args.scale_s, args.cache_batch_idx,
args.scale_p, args.scale_s,
args.stride_q, args.scale_p,
args.stride_k, args.stride_q,
args.stride_v, args.stride_k,
args.stride_bias, args.stride_v,
(1 < args.num_splits ? args.stride_o_acc : args.stride_o), args.stride_bias,
args.nhead_stride_q, (1 < args.num_splits ? args.stride_o_acc : args.stride_o),
args.nhead_stride_k, args.nhead_stride_q,
args.nhead_stride_v, args.nhead_stride_k,
args.nhead_stride_bias, args.nhead_stride_v,
(1 < args.num_splits ? args.nhead_stride_lse_acc : args.nhead_stride_lse), args.nhead_stride_bias,
(1 < args.num_splits ? args.nhead_stride_o_acc : args.nhead_stride_o), (1 < args.num_splits ? args.nhead_stride_lse_acc : args.nhead_stride_lse),
args.batch_stride_q, (1 < args.num_splits ? args.nhead_stride_o_acc : args.nhead_stride_o),
args.batch_stride_k, args.batch_stride_q,
args.batch_stride_v, args.batch_stride_k,
args.batch_stride_bias, args.batch_stride_v,
(1 < args.num_splits ? args.batch_stride_lse_acc : args.batch_stride_lse), args.batch_stride_bias,
(1 < args.num_splits ? args.batch_stride_o_acc : args.batch_stride_o), (1 < args.num_splits ? args.batch_stride_lse_acc : args.batch_stride_lse),
(1 < args.num_splits ? args.split_stride_lse_acc : 0), (1 < args.num_splits ? args.batch_stride_o_acc : args.batch_stride_o),
(1 < args.num_splits ? args.split_stride_o_acc : 0), (1 < args.num_splits ? args.split_stride_lse_acc : 0),
args.window_size_left, (1 < args.num_splits ? args.split_stride_o_acc : 0),
args.window_size_right, args.window_size_left,
args.mask_type); args.window_size_right,
args.mask_type);
} }
}(); }();
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment