diff --git a/example/ck_tile/36_quant/example_dynamic_per_tensor_quant.cpp b/example/ck_tile/36_quant/example_dynamic_per_tensor_quant.cpp index 06a9247e31..15d00995c5 100644 --- a/example/ck_tile/36_quant/example_dynamic_per_tensor_quant.cpp +++ b/example/ck_tile/36_quant/example_dynamic_per_tensor_quant.cpp @@ -36,20 +36,20 @@ auto get_elimit() int main() { - static constexpr ck_tile::index_t Repeat_M_ = 1; + static constexpr ck_tile::index_t Repeat_M_ = 8; static constexpr ck_tile::index_t Repeat_N_ = 1; - static constexpr ck_tile::index_t ThreadPerBlock_M_ = 4; + static constexpr ck_tile::index_t ThreadPerBlock_M_ = 8; static constexpr ck_tile::index_t ThreadPerBlock_N_ = 64; - static constexpr ck_tile::index_t Vector_N_ = 2; + static constexpr ck_tile::index_t Vector_N_ = 1; static constexpr bool is_warp_per_row = ThreadPerBlock_N_ <= warpSize; static_assert((ThreadPerBlock_M_ * ThreadPerBlock_N_) % warpSize == 0); static constexpr ck_tile::index_t total_warps = (ThreadPerBlock_M_ * ThreadPerBlock_N_) / warpSize; - - // num of warps along mreference_static_per_tensor_quantization2d + // std::cout<<"total_warps: "<; - int m = 256; - int n = 256; - int x_stride = 256; + int m = 64; + int n = 64; + int x_stride = 64; ck_tile::HostTensor x_host({m, n}, {x_stride, 1}); ck_tile::HostTensor scale_host({1}, {1}); @@ -144,7 +144,7 @@ int main() scale_buf.FromDevice(scale_host.data()); ck_tile::reference_per_tensor_quantization2d( x_host, scale_host, qx_host_ref); - + // std::cout<(); diff --git a/include/ck_tile/ops/quant/pipeline/quant_pipeline.hpp b/include/ck_tile/ops/quant/pipeline/quant_pipeline.hpp index fac15254ff..756282623a 100644 --- a/include/ck_tile/ops/quant/pipeline/quant_pipeline.hpp +++ b/include/ck_tile/ops/quant/pipeline/quant_pipeline.hpp @@ -9,7 +9,6 @@ #include namespace ck_tile { - template struct StaticPerTensorQuantPipeline { @@ -92,7 +91,7 @@ struct DynamicPerTensorQuantPipeline { auto x_window = make_tile_window(x_window_, Policy::template MakeXBlockTileDistribution()); - + auto origin = x_window.get_window_origin(); static constexpr index_t Block_N = Problem::BlockShape::Block_N; index_t num_n_tile_iteration = __builtin_amdgcn_readfirstlane(integer_divide_ceil(row_size, Block_N)); @@ -130,12 +129,8 @@ struct DynamicPerTensorQuantPipeline block_reduce2d_cross_warp_sync(absmax, smem, reduce_max_func); *scale = absmax.get_thread_buffer()[0] / ck_tile::numeric::max(); - ck_tile::index_t stride_to_right_most_window = - row_size % Block_N == 0 ? row_size - Block_N : row_size - row_size % Block_N; - move_tile_window(x_window, {0, -Block_N}); - move_tile_window(qx_window, {0, stride_to_right_most_window}); - for(int iN = __builtin_amdgcn_readfirstlane(0); iN < num_n_tile_iteration; ++iN) - { + x_window.set_window_origin(origin); + for(int iN = __builtin_amdgcn_readfirstlane(0); iN < num_n_tile_iteration; ++iN){ const auto x = load_tile(x_window); const auto qx = tile_elementwise_in( [&](const auto& a) { @@ -143,8 +138,8 @@ struct DynamicPerTensorQuantPipeline }, x); store_tile(qx_window, qx); - move_tile_window(x_window, {0, -Block_N}); - move_tile_window(qx_window, {0, -Block_N}); + move_tile_window(x_window, {0, Block_N}); + move_tile_window(qx_window, {0, Block_N}); } } };