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

Remove debug macro usages

Browse Source
This commit is contained in:
PoYen, Chen
2024-08-07 15:26:43 +00:00
parent b6c2f2f01d
commit cef9da0a76
2 changed files with 3 additions and 54 deletions
Download Patch File Download Diff File Expand all files Collapse all files

46
include/ck_tile/ops/fmha/kernel/fmha_fwd_appendkv_kernel.hpp
View File

@@ -403,16 +403,6 @@ struct FmhaFwdAppendKVKernel
const index_t num_blocks =
integer_divide_ceil(kargs.seqlen_k + kargs.seqlen_knew, kargs.page_block_size);
DEVICE_DEBUG_STMTS
{
printf("[DEVICE] block_indics: ");
for(index_t i_block = 0; i_block < num_blocks; ++i_block)
{
printf("(%d, %d) ", i_block, block_indices[i_block]);
}
printf("\n");
}
const long_index_t fixed_offset =
static_cast<long_index_t>(i_nhead_ / kargs.nhead_ratio_qk) *
kargs.nhead_stride_v;
@@ -723,14 +713,7 @@ struct FmhaFwdAppendKVKernel
auto [i_block0, k_dram_window_tmp] = k_tile_navigator.make_tile_window(
k_dram_window, {skip_append_kv ? 0 : kargs.seqlen_k + i_n0, 0});
DEVICE_DEBUG_STMTS
{
printf("[DEVICE] i_block0: %d\n", i_block0);
auto local_origin = k_dram_window_tmp.get_window_origin();
printf("[DEVICE] origin: (%d, %d)\n",
local_origin.at(number<0>{}),
local_origin.at(number<1>{}));
}
auto knew_dram_window =
make_tile_window(knew_dram,
make_tuple(number<FmhaPipeline::kN0>{}, number<FmhaPipeline::kK0>{}),
@@ -744,37 +727,12 @@ struct FmhaFwdAppendKVKernel
auto [i_block1, v_dram_window_tmp] = v_tile_navigator.make_tile_window(
v_dram_window, {0, skip_append_kv ? 0 : kargs.seqlen_k + i_n0});
if constexpr(kIsPagedKV)
{
DEVICE_DEBUG_STMTS
{
printf("[DEVICE] i_block1: %d\n", i_block1);
auto local_origin = v_dram_window_tmp.get_window_origin();
printf("[DEVICE] origin: (%d, %d)\n",
local_origin.at(number<0>{}),
local_origin.at(number<1>{}));
printf("[DEVICE] psychical block_ptr 0: %p\n",
static_cast<void*>(v_tile_navigator.physical_blocks +
0 * v_tile_navigator.block_stride));
printf("[DEVICE] psychical block_ptr 1: %p\n",
static_cast<void*>(v_tile_navigator.physical_blocks +
1 * v_tile_navigator.block_stride));
printf("[DEVICE] tile window data ptr: %p\n",
static_cast<void*>(v_dram_window_tmp.get_bottom_tensor_view().buf_.p_data_));
}
}
auto vnew_dram_window =
make_tile_window(vnew_dram,
make_tuple(number<FmhaPipeline::kN1>{}, number<FmhaPipeline::kN0>{}),
{0, i_n0});
DEVICE_DEBUG_STMTS
{
printf("[DEVICE] skip_transform_q: %d, skip_appendkv: %d\n",
kargs.seqlen_q <= i_m0,
kargs.seqlen_knew <= i_n0);
}
if constexpr(kApplyRoPE)
{
FmhaPipeline{}(q_dram_window,

11
include/ck_tile/ops/fmha/pipeline/block_fmha_fwd_splitkv_pipeline_qr_ks_vs.hpp
View File

@@ -231,16 +231,7 @@ struct BlockFmhaFwdSplitKVPipelineQRKSVS
}();
const auto num_total_loop =
integer_divide_ceil(seqlen_k_end - adjusted_seqlen_k_start, kN0);
#if 0
DEVICE_DEBUG_STMTS
{
printf("[DEVICE] seqlen_k_start: %d, seqlen_k_end: %d\n", seqlen_k_start, seqlen_k_end);
printf("[DEVICE] adjusted_seqlen_k_start: %d, num_total_loop: %d\n",
adjusted_seqlen_k_start,
num_total_loop);
printf("[DEVICE] kHasUnevenSplits: %d\n", kHasUnevenSplits);
}
#endif
// check early exit if masked and no work to do.
if constexpr(FmhaMask::IsMasking || kHasUnevenSplits)
{