diff --git a/include/ck_tile/core/tensor/tile_scatter_gather.hpp b/include/ck_tile/core/tensor/tile_scatter_gather.hpp index 9d16b7cde2..c8c7e5eb6d 100644 --- a/include/ck_tile/core/tensor/tile_scatter_gather.hpp +++ b/include/ck_tile/core/tensor/tile_scatter_gather.hpp @@ -522,18 +522,22 @@ struct tile_scatter_gather constexpr auto idx_gather = idx_ys_start[number{}]; const auto page_offset = page_idx_[idx_gather]; + // merge page_offset into bottom_coord + auto mixed_bottom_thread_coord = bottom_tensor_thread_coord; + mixed_bottom_thread_coord.get_hidden_index()[number<0>{}] += page_offset; + // read from bottom tensor if constexpr(std::is_same_v) this->get_bottom_tensor_view().template async_get_vectorized_elements( smem, - bottom_tensor_thread_coord, - page_offset, + mixed_bottom_thread_coord, + number<0>{}, bool_constant{}); else this->get_bottom_tensor_view().template async_get_vectorized_elements( smem, - bottom_tensor_thread_coord, - page_offset, + mixed_bottom_thread_coord, + number<0>{}, valids_[idx_gather], bool_constant{}); diff --git a/include/ck_tile/ops/flatmm/pipeline/mixed_prec_flatmm_pipeline_agmem_bgmem_creg_v1_policy.hpp b/include/ck_tile/ops/flatmm/pipeline/mixed_prec_flatmm_pipeline_agmem_bgmem_creg_v1_policy.hpp index fb2e4de0ce..01535d06dd 100644 --- a/include/ck_tile/ops/flatmm/pipeline/mixed_prec_flatmm_pipeline_agmem_bgmem_creg_v1_policy.hpp +++ b/include/ck_tile/ops/flatmm/pipeline/mixed_prec_flatmm_pipeline_agmem_bgmem_creg_v1_policy.hpp @@ -7,7 +7,7 @@ namespace ck_tile { -#define CKTILE_FLATMM_USE_BUFFER_LOAD_LDS_AS_POSSIBLE 1 +#define CKTILE_FLATMM_USE_BUFFER_LOAD_LDS_AS_POSSIBLE 0 #if defined(__gfx950__) #define CKTILE_FLATMM_ARCH_SUPPORT_BUFFER_LOAD_LDS_DWORDx4 1