From 19c19d8bd3f0a73bbf6b38c639091414ae4fbcd4 Mon Sep 17 00:00:00 2001 From: "PoYen, Chen" Date: Tue, 13 Aug 2024 07:26:26 +0000 Subject: [PATCH] Only expose necessary methods (not attributes) --- .../ops/fmha/block/page_block_navigator.hpp | 30 ++++++++----------- 1 file changed, 12 insertions(+), 18 deletions(-) diff --git a/include/ck_tile/ops/fmha/block/page_block_navigator.hpp b/include/ck_tile/ops/fmha/block/page_block_navigator.hpp index 6b77e9cc19..82e93c128c 100644 --- a/include/ck_tile/ops/fmha/block/page_block_navigator.hpp +++ b/include/ck_tile/ops/fmha/block/page_block_navigator.hpp @@ -46,13 +46,6 @@ struct TrivialPageBlockNavigator return /*block_index=*/0; } - /// TODO: remove this method after finish debuging - CK_TILE_DEVICE static constexpr int32_t - get_block_index(const WindowOrigin& /*global_window_origin*/) - { - return /*block_index=*/0; - } - CK_TILE_DEVICE static constexpr WindowOrigin to_local_window_origin(const WindowOrigin& global_window_origin) { @@ -169,17 +162,6 @@ struct PageBlockNavigator tile_window.set_bottom_tensor_view_data_ptr(get_block_ptr(new_block_index)); } - CK_TILE_HOST_DEVICE - DataType* get_block_ptr(index_t block_index) const - { - return physical_blocks + physical_block_indices[block_index] * block_stride + fixed_offset; - } - - CK_TILE_HOST_DEVICE int32_t get_block_index(const WindowOrigin& global_window_origin) const - { - return integer_divide_floor(global_window_origin.at(number{}), page_block_size); - } - CK_TILE_HOST_DEVICE WindowOrigin to_local_window_origin(const WindowOrigin& global_window_origin) const { @@ -216,6 +198,18 @@ struct PageBlockNavigator } } + private: + CK_TILE_HOST_DEVICE + DataType* get_block_ptr(index_t block_index) const + { + return physical_blocks + physical_block_indices[block_index] * block_stride + fixed_offset; + } + + CK_TILE_HOST_DEVICE int32_t get_block_index(const WindowOrigin& global_window_origin) const + { + return integer_divide_floor(global_window_origin.at(number{}), page_block_size); + } + DataType* physical_blocks; long_index_t block_stride; long_index_t fixed_offset;