From 534d4594d0906288339c937b6419f5995ef09889 Mon Sep 17 00:00:00 2001 From: Aviral Goel Date: Thu, 22 May 2025 01:28:00 -0500 Subject: [PATCH 1/5] Refactor tile_window.hpp, tile_window_linear.hpp into a CK Tile Hierarchy (#2214) * window_origin variable now in base class * abstracted more functions * consolidated tile_window_static_distribution and tile_window_static_lengths * clang format * skeleton code for tile_window and tile_window_linear consolidation * more abstraction * moved variables from child to parent * clang format * removed comments * removed debug code * removed debug code * abstracting traits WIP * consolidated traits * removed comments and clang formatted --- include/ck_tile/core.hpp | 1 + include/ck_tile/core/tensor/tile_window.hpp | 540 +++++------------ .../ck_tile/core/tensor/tile_window_base.hpp | 256 +++++++++ .../core/tensor/tile_window_linear.hpp | 544 ++++++------------ 4 files changed, 571 insertions(+), 770 deletions(-) create mode 100644 include/ck_tile/core/tensor/tile_window_base.hpp diff --git a/include/ck_tile/core.hpp b/include/ck_tile/core.hpp index 2ea8bf15a7..aa9411b2e1 100644 --- a/include/ck_tile/core.hpp +++ b/include/ck_tile/core.hpp @@ -54,6 +54,7 @@ #include "ck_tile/core/tensor/tile_distribution_encoding.hpp" #include "ck_tile/core/tensor/tile_elementwise.hpp" #include "ck_tile/core/tensor/tile_scatter_gather.hpp" +#include "ck_tile/core/tensor/tile_window_base.hpp" #include "ck_tile/core/tensor/tile_window.hpp" #include "ck_tile/core/tensor/tile_window_linear.hpp" #include "ck_tile/core/tensor/tile_window_utils.hpp" diff --git a/include/ck_tile/core/tensor/tile_window.hpp b/include/ck_tile/core/tensor/tile_window.hpp index 716b1f4ecb..d8a5c14f9b 100644 --- a/include/ck_tile/core/tensor/tile_window.hpp +++ b/include/ck_tile/core/tensor/tile_window.hpp @@ -13,6 +13,7 @@ #include "ck_tile/core/tensor/static_distributed_tensor.hpp" #include "ck_tile/core/tensor/tensor_adaptor.hpp" #include "ck_tile/core/tensor/tile_distribution.hpp" +#include "ck_tile/core/tensor/tile_window_base.hpp" #include "ck_tile/core/utility/functional.hpp" #include "ck_tile/core/utility/type_traits.hpp" @@ -34,166 +35,60 @@ template struct tile_window_with_static_distribution + : public tile_window_with_tile_dstr_base< + tile_window_with_static_distribution, + BottomTensorView_, + WindowLengths_, + StaticTileDistribution_> { - using BottomTensorView = remove_reference_t; - using WindowLengths = remove_cvref_t; - using TileDstr = remove_cvref_t; - - using WindowAdaptor = typename TileDstr::PsYs2XsAdaptor; - using BottomTensorDesc = typename BottomTensorView::TensorDesc; - - using DataType = remove_cvref_t; - - static constexpr index_t NDimWindowAdaptorTop = WindowAdaptor::get_num_of_top_dimension(); - static constexpr index_t NDimBottomTensor = BottomTensorDesc::get_num_of_dimension(); - - static constexpr index_t NDimP = TileDstr::get_num_of_dimension_p(); - static constexpr index_t NDimY = TileDstr::get_num_of_dimension_y(); + using Base = tile_window_with_tile_dstr_base< + tile_window_with_static_distribution, + BottomTensorView_, + WindowLengths_, + StaticTileDistribution_>; static constexpr auto I0 = number<0>{}; static constexpr auto I1 = number<1>{}; static_assert(NumCoord == 1); - // TODO: check WindowLengths and StaticTileDistribution are consistent - - static_assert(ck_tile::is_known_at_compile_time::value, - "wrong! lengths should be static"); - static_assert(TileDstr::is_static(), "wrong!"); - - static_assert(NDimBottomTensor == WindowAdaptor::get_num_of_bottom_dimension(), - "wrong! inconsistent # of diemsnions"); - - using AdaptorTopIndex = array; - using BottomTensorIndex = array; - - using WindowAdaptorCoord = - decltype(make_tensor_adaptor_coordinate(WindowAdaptor{}, AdaptorTopIndex{})); - - using BottomTensorCoord = - decltype(make_tensor_coordinate(BottomTensorDesc{}, BottomTensorIndex{})); - - struct load_store_traits - { - private: - static constexpr auto get_vector_dim_y_scalar_per_vector() - { - const auto [ys_vector_lengths, ys_vector_strides] = - tile_window_with_static_distribution:: - get_window_adaptor_ys_safe_vector_length_strides(); - - index_t VectorDimY_ = 0; - index_t ScalarPerVector_ = 1; - - for(index_t i = 0; i < NDimY; ++i) - { - if(ys_vector_strides[i] == 1 && ys_vector_lengths[i] > ScalarPerVector_) - { - ScalarPerVector_ = ys_vector_lengths[i]; - VectorDimY_ = i; - } - } - - return make_tuple(VectorDimY_, ScalarPerVector_); - } - - public: - static constexpr index_t PackedSize = - ck_tile::numeric_traits>::PackedSize; - static constexpr index_t VectorDimY = get_vector_dim_y_scalar_per_vector().template at<0>(); - static constexpr index_t ScalarPerVector = - get_vector_dim_y_scalar_per_vector().template at<1>(); - - // using vector_type_t = vector_type_maker_t; - // using vector_t = typename vector_type_t::type; - using vector_t = thread_buffer; - - private: - static constexpr auto scalars_per_access_ = [] { - constexpr auto scalars_per_access_arr = generate_array( - [&](auto i) { return (i == VectorDimY) ? ScalarPerVector : 1; }, number{}); - - /// TODO: add non-automatic storage argument support to macro TO_SEQUENCE() - constexpr auto NDimY_ = NDimY; - - return TO_SEQUENCE(scalars_per_access_arr, NDimY_); - }(); - - static constexpr auto get_space_filling_curve() - { - constexpr auto tile_dstr = TileDstr{}; - - constexpr auto thread_tensor_lengths_ys = - to_sequence(tile_dstr.get_ys_to_d_descriptor().get_lengths()); - - // FIXME: need logic to judge dim access order - using DimAccessOrder = typename arithmetic_sequence_gen<0, NDimY, 1>::type; - - return space_filling_curve{}; - } - - public: - using SFC_Ys = decltype(get_space_filling_curve()); - - static constexpr index_t NumAccess = SFC_Ys::get_num_of_access(); - - static_assert(0 < NumAccess, "Wrong! NumAccess should be larger than 0"); - static_assert(NumAccess % NumCoord == 0, "wrong! # of access is not divisible by NumCoord"); - }; - - static constexpr index_t NumAccessPerCoord = load_store_traits::NumAccess / NumCoord; + static_assert(Base::Traits::NumAccess % NumCoord == 0, + "wrong! # of access is not divisible by NumCoord"); + static constexpr index_t NumAccessPerCoord = Base::Traits::NumAccess / NumCoord; CK_TILE_DEVICE constexpr tile_window_with_static_distribution() = default; CK_TILE_DEVICE constexpr tile_window_with_static_distribution( - const BottomTensorView& bottom_tensor_view, - const WindowLengths& window_lengths, - const BottomTensorIndex& window_origin, - const TileDstr& tile_distribution) - : bottom_tensor_view_{bottom_tensor_view}, - window_lengths_{window_lengths}, - window_origin_{window_origin}, - tile_dstr_{tile_distribution}, - pre_computed_coords_{} + const typename Base::BottomTensorView& bottom_tensor_view, + const typename Base::WindowLengths& window_lengths, + const typename Base::BottomTensorIndex& window_origin, + const typename Base::TileDstr& tile_distribution) + : pre_computed_coords_{} { -#if 0 // debug - // TODO: this use more register for FA, but less register for GEMM - // need investigation - // only support warp-tile and block-tile - static_assert(NDimP == 1 or NDimP == 2, "wrong!"); - WindowAdaptorCoord window_adaptor_thread_coord_tmp; - - if constexpr(NDimP == 1) - { - window_adaptor_thread_coord_tmp = make_tensor_adaptor_coordinate( - tile_distribution.get_ps_ys_to_xs_adaptor(), AdaptorTopIndex{get_lane_id(), 0}); - } - else if constexpr(NDimP == 2) - { - window_adaptor_thread_coord_tmp = - make_tensor_adaptor_coordinate(tile_distribution.get_ps_ys_to_xs_adaptor(), - AdaptorTopIndex{get_warp_id(), get_lane_id(), 0}); - } -#else - // TODO: this use less register for FA, but more register for GEMM - // need investigation + this->window_origin_ = window_origin; + this->window_lengths_ = window_lengths; + this->bottom_tensor_view_ = bottom_tensor_view; + this->tile_dstr_ = tile_distribution; const auto window_adaptor_thread_coord_tmp = make_tensor_adaptor_coordinate( tile_distribution.get_ps_ys_to_xs_adaptor(), container_concat(detail::get_partition_index(tile_distribution), - array{0})); -#endif + array{0})); - BottomTensorIndex bottom_tensor_thread_origin_idx_tmp = + typename Base::BottomTensorIndex bottom_tensor_thread_origin_idx_tmp = window_origin + window_adaptor_thread_coord_tmp.get_bottom_index(); const auto bottom_tensor_thread_coord_tmp = make_tensor_coordinate( - bottom_tensor_view_.get_tensor_descriptor(), bottom_tensor_thread_origin_idx_tmp); + bottom_tensor_view.get_tensor_descriptor(), bottom_tensor_thread_origin_idx_tmp); // pre-compute NumCoord (WindowAdaptorCoord, BottomTensorCoord) bundles to speed up // future load/store() calls (might allocate more registers) - using Traits = load_store_traits; + using Traits = typename Base::Traits; using SFC_Ys = typename Traits::SFC_Ys; static_for<0, NumCoord, 1>{}([&](auto iCoord) { @@ -204,9 +99,10 @@ struct tile_window_with_static_distribution SFC_Ys::get_step_between(number<0>{}, number{}); constexpr auto idx_diff_ps_ys = container_concat( - generate_tuple([&](auto) { return number<0>{}; }, number{}), idx_diff_ys); + generate_tuple([&](auto) { return number<0>{}; }, number{}), + idx_diff_ys); - move_window_adaptor_and_bottom_tensor_thread_coordinate( + Base::move_window_adaptor_and_bottom_tensor_thread_coordinate( window_adaptor_thread_coord, bottom_tensor_thread_coord, idx_diff_ps_ys); pre_computed_coords_(iCoord) = @@ -214,95 +110,12 @@ struct tile_window_with_static_distribution }); } - CK_TILE_DEVICE static constexpr index_t get_num_of_dimension() { return NDimBottomTensor; } - - CK_TILE_DEVICE static constexpr bool has_static_tile_distribution() - { - return TileDstr::is_static(); - } - - CK_TILE_DEVICE constexpr auto get_window_lengths() const { return window_lengths_; } - - CK_TILE_DEVICE constexpr auto get_tile_distribution() const { return tile_dstr_; } - - CK_TILE_DEVICE constexpr auto get_bottom_tensor_view() const { return bottom_tensor_view_; } - - CK_TILE_DEVICE constexpr auto get_window_origin() const { return window_origin_; } - - CK_TILE_DEVICE constexpr void - set_bottom_tensor_view_data_ptr(typename BottomTensorView::DataType* data) - { - bottom_tensor_view_.buf_.p_data_ = data; - } - - // move thread's window adaptor coordinate and bottom tensor coordinate - // [p0, p1, ..., y0, y1, ...] ==> [x0, x1, ...] ==> [x0', x1', ...] ==> [offset] - template - CK_TILE_DEVICE void move_window_adaptor_and_bottom_tensor_thread_coordinate( - WindowAdaptorCoord& window_adaptor_thread_coord, - BottomTensorCoord& bottom_tensor_thread_coord, - const ATopIndex& idx_diff_adaptor_top) const - { - array idx_diff_adaptor_bottom; - - move_tensor_adaptor_coordinate(tile_dstr_.get_ps_ys_to_xs_adaptor(), - window_adaptor_thread_coord, - idx_diff_adaptor_top, - idx_diff_adaptor_bottom); - - move_tensor_coordinate(bottom_tensor_view_.get_tensor_descriptor(), - bottom_tensor_thread_coord, - idx_diff_adaptor_bottom); - } - - // return vector dimension among [y0, y1, ...] - CK_TILE_DEVICE static constexpr auto get_window_adaptor_ys_safe_vector_length_strides() - { - // bottom tensor top dimension vector lengths and strides - const auto [bottom_tensor_top_dim_vector_lengths, bottom_tensor_top_dim_vector_strides] = - BottomTensorDesc::get_top_dimension_safe_vector_length_strides(); - - // window vector lengths/strides - const auto window_adaptor_bottom_dim_vector_lengths = bottom_tensor_top_dim_vector_lengths; - const auto window_adaptor_bottom_dim_vector_strides = bottom_tensor_top_dim_vector_strides; - - // window adaptor [p0, p1, ..., y0, y1, ...] - array window_adaptor_vector_lengths{ - -1}; - array window_adaptor_vector_strides{ - -1}; - - constexpr auto window_adaptor_bottom_dims = - WindowAdaptor::get_bottom_dimension_hidden_ids(); - - set_container_subset(window_adaptor_vector_lengths, - window_adaptor_bottom_dims, - window_adaptor_bottom_dim_vector_lengths); - set_container_subset(window_adaptor_vector_strides, - window_adaptor_bottom_dims, - window_adaptor_bottom_dim_vector_strides); - - const auto [window_adaptor_ps_ys_vector_lengths, window_adaptor_ps_ys_vector_strides] = - WindowAdaptor{}.get_top_dimension_safe_vector_length_strides( - window_adaptor_vector_lengths, window_adaptor_vector_strides); - - // [y0, y1, ...] - constexpr auto y_dims = typename arithmetic_sequence_gen::type{}; - - return make_tuple(get_container_subset(window_adaptor_ps_ys_vector_lengths, y_dims), - get_container_subset(window_adaptor_ps_ys_vector_strides, y_dims)); - } - - CK_TILE_DEVICE constexpr auto get_num_of_access() const { return load_store_traits::NumAccess; } - template CK_TILE_DEVICE auto load(number = {}, bool_constant = {}) const { - constexpr auto tile_dstr = TileDstr{}; - auto dst_tensor = make_static_distributed_tensor(tile_dstr); + constexpr auto tile_dstr = typename Base::TileDstr{}; + auto dst_tensor = make_static_distributed_tensor(tile_dstr); load(dst_tensor, number{}, bool_constant{}); return dst_tensor; } @@ -314,11 +127,11 @@ struct tile_window_with_static_distribution number = {}, bool_constant = {}) const { - using Traits = load_store_traits; + using Traits = typename Base::Traits; using vector_t = typename Traits::vector_t; using SFC_Ys = typename Traits::SFC_Ys; - constexpr auto tile_dstr = TileDstr{}; + constexpr auto tile_dstr = typename Base::TileDstr{}; // loop over thread tensor space [y0, y1, ...] static_for<0, NumCoord, 1>{}([&](auto iCoord) { @@ -334,9 +147,8 @@ struct tile_window_with_static_distribution // read from bottom tensor const vector_t vec_value = - get_bottom_tensor_view().template get_vectorized_elements( + this->get_bottom_tensor_view().template get_vectorized_elements( bottom_tensor_thread_coord, 0, bool_constant{}); -#if 1 // write into distributed tensor static_for<0, Traits::ScalarPerVector, Traits::PackedSize>{}([&](auto j) { constexpr auto idx_ys = generate_tuple( @@ -344,33 +156,26 @@ struct tile_window_with_static_distribution return jj == Traits::VectorDimY ? (idx_ys_start[jj] + j) : idx_ys_start[jj]; }, - number{}); + number{}); constexpr index_t d = tile_dstr.get_ys_to_d_descriptor().calculate_offset(idx_ys) / Traits::PackedSize; dst_tensor.get_thread_buffer().template at() = - vec_value.template get_as()[j / Traits::PackedSize]; + vec_value + .template get_as()[j / Traits::PackedSize]; }); -#else - constexpr index_t d = - tile_dstr.get_ys_to_d_descriptor().calculate_offset(idx_ys_start); - static_assert(d % Traits::ScalarPerVector == 0); - - dst_tensor.get_thread_buffer().template get_as()( - number{}) = bit_cast(vec_value); -#endif // move thread coordinate if constexpr(iCoordAccess != (NumAccessPerCoord - 1)) { constexpr auto idx_diff_ys = SFC_Ys::get_forward_step(iAccess); constexpr auto idx_diff_ps_ys = container_concat( - generate_tuple([&](auto) { return number<0>{}; }, number{}), + generate_tuple([&](auto) { return number<0>{}; }, number{}), idx_diff_ys); - move_window_adaptor_and_bottom_tensor_thread_coordinate( + Base::move_window_adaptor_and_bottom_tensor_thread_coordinate( window_adaptor_thread_coord, bottom_tensor_thread_coord, idx_diff_ps_ys); } }); @@ -386,22 +191,16 @@ struct tile_window_with_static_distribution bool_constant = {}, bool_constant = {}) const { - using Traits = load_store_traits; - - // using vector_type_t = typename Traits::vector_type_t; + using Traits = typename Base::Traits; using vector_t = typename Traits::vector_t; using SFC_Ys = typename Traits::SFC_Ys; static constexpr index_t YElementSize = - TileDstr{}.get_ys_to_d_descriptor().get_element_space_size(); + typename Base::TileDstr{}.get_ys_to_d_descriptor().get_element_space_size(); static_assert(YElementSize % (Traits::PackedSize * Traits::ScalarPerVector) == 0); using vectorized_tbuf = array; - // StaticBuffer; - constexpr auto tile_dstr = TileDstr{}; + constexpr auto tile_dstr = typename Base::TileDstr{}; auto& dst_vec_tbuf = reinterpret_cast(dst_tensor.get_thread_buffer()); @@ -427,7 +226,7 @@ struct tile_window_with_static_distribution Traits::PackedSize; static_assert(d % Traits::ScalarPerVector == 0); - get_bottom_tensor_view().template get_vectorized_elements_raw( + this->get_bottom_tensor_view().template get_vectorized_elements_raw( dst_vec_tbuf.template at(), bottom_tensor_thread_coord, 0 /**/, @@ -444,10 +243,10 @@ struct tile_window_with_static_distribution constexpr auto idx_diff_ys = SFC_Ys::get_forward_step(iAccess); constexpr auto idx_diff_ps_ys = container_concat( - generate_tuple([&](auto) { return number<0>{}; }, number{}), + generate_tuple([&](auto) { return number<0>{}; }, number{}), idx_diff_ys); - move_window_adaptor_and_bottom_tensor_thread_coordinate( + Base::move_window_adaptor_and_bottom_tensor_thread_coordinate( window_adaptor_thread_coord, bottom_tensor_thread_coord, idx_diff_ps_ys); } }); @@ -492,9 +291,8 @@ struct tile_window_with_static_distribution const index_t m0_init_value = size_per_buf + size_per_wave * get_warp_id(); m0_set_with_memory(m0_init_value); // This should be wave independent - using Traits = load_store_traits; + using Traits = typename Base::Traits; - // using vector_type_t = typename Traits::vector_type_t; using vector_t = typename Traits::vector_t; using SFC_Ys = typename Traits::SFC_Ys; @@ -516,7 +314,7 @@ struct tile_window_with_static_distribution }(); // read from bottom tensor - get_bottom_tensor_view().template async_get_vectorized_elements_raw( + this->get_bottom_tensor_view().template async_get_vectorized_elements_raw( smem, bottom_tensor_thread_coord, 0, pre_nop_); // move thread coordinate @@ -525,10 +323,10 @@ struct tile_window_with_static_distribution constexpr auto idx_diff_ys = SFC_Ys::get_forward_step(iAccess); constexpr auto idx_diff_ps_ys = container_concat( - generate_tuple([&](auto) { return number<0>{}; }, number{}), + generate_tuple([&](auto) { return number<0>{}; }, number{}), idx_diff_ys); - move_window_adaptor_and_bottom_tensor_thread_coordinate( + Base::move_window_adaptor_and_bottom_tensor_thread_coordinate( window_adaptor_thread_coord, bottom_tensor_thread_coord, idx_diff_ps_ys); m0_inc_with_memory(size_per_issue); @@ -569,7 +367,7 @@ struct tile_window_with_static_distribution const index_t m0_init_value = size_per_buf + size_per_wave * get_warp_id(); - using Traits = load_store_traits; + using Traits = typename Base::Traits; using vector_t = typename Traits::vector_t; using SFC_Ys = typename Traits::SFC_Ys; @@ -588,7 +386,7 @@ struct tile_window_with_static_distribution constexpr auto iAccess = number{}; // read from bottom tensor - get_bottom_tensor_view().template async_get_vectorized_elements( + this->get_bottom_tensor_view().template async_get_vectorized_elements( smem, bottom_tensor_thread_coord, 0, bool_constant{}); // move thread coordinate @@ -597,10 +395,10 @@ struct tile_window_with_static_distribution constexpr auto idx_diff_ys = SFC_Ys::get_forward_step(iAccess); constexpr auto idx_diff_ps_ys = container_concat( - generate_tuple([&](auto) { return number<0>{}; }, number{}), + generate_tuple([&](auto) { return number<0>{}; }, number{}), idx_diff_ys); - move_window_adaptor_and_bottom_tensor_thread_coordinate( + Base::move_window_adaptor_and_bottom_tensor_thread_coordinate( window_adaptor_thread_coord, bottom_tensor_thread_coord, idx_diff_ps_ys); smem += size_per_issue; // Note we manually increase the per-issue offset @@ -610,17 +408,18 @@ struct tile_window_with_static_distribution } template - CK_TILE_DEVICE void store(const static_distributed_tensor& dstr_tensor, + CK_TILE_DEVICE void store(const static_distributed_tensor& dstr_tensor, number = {}, bool_constant = {}) const { - using Traits = load_store_traits; + using Traits = typename Base::Traits; // using vector_type_t = typename Traits::vector_type_t; using vector_t = typename Traits::vector_t; using SFC_Ys = typename Traits::SFC_Ys; - constexpr auto tile_dstr = TileDstr{}; + constexpr auto tile_dstr = typename Base::TileDstr{}; // loop over thread tensor space [y0, y1, ...] static_for<0, NumCoord, 1>{}([&](auto iCoord) { @@ -643,20 +442,20 @@ struct tile_window_with_static_distribution return jj == Traits::VectorDimY ? (idx_ys_start[jj] + j) : idx_ys_start[jj]; }, - number{}); + number{}); constexpr index_t d = tile_dstr.get_ys_to_d_descriptor().calculate_offset(idx_ys) / Traits::PackedSize; - vec_value.template get_as()(j / Traits::PackedSize) = + vec_value.template get_as()(j / Traits::PackedSize) = dstr_tensor.get_thread_buffer().template at(); }); // const vector_t vec_value = vec.template get_as().template at<0>(); // write into bottom tensor - get_bottom_tensor_view().template set_vectorized_elements( + this->get_bottom_tensor_view().template set_vectorized_elements( bottom_tensor_thread_coord, 0, vec_value, @@ -668,10 +467,10 @@ struct tile_window_with_static_distribution constexpr auto idx_diff_ys = SFC_Ys::get_forward_step(iAccess); constexpr auto idx_diff_ps_ys = container_concat( - generate_tuple([&](auto) { return number<0>{}; }, number{}), + generate_tuple([&](auto) { return number<0>{}; }, number{}), idx_diff_ys); - move_window_adaptor_and_bottom_tensor_thread_coordinate( + Base::move_window_adaptor_and_bottom_tensor_thread_coordinate( window_adaptor_thread_coord, bottom_tensor_thread_coord, idx_diff_ps_ys); } }); @@ -679,15 +478,17 @@ struct tile_window_with_static_distribution } template - CK_TILE_DEVICE void store_raw(const static_distributed_tensor& dstr_tensor, - number = {}) const + CK_TILE_DEVICE void + store_raw(const static_distributed_tensor& + dstr_tensor, + number = {}) const { - using Traits = load_store_traits; + using Traits = typename Base::Traits; using vector_t = typename Traits::vector_t; using SFC_Ys = typename Traits::SFC_Ys; - constexpr auto tile_dstr = TileDstr{}; + constexpr auto tile_dstr = typename Base::TileDstr{}; static constexpr bool oob_conditional_check = true; // loop over thread tensor space [y0, y1, ...] @@ -710,16 +511,16 @@ struct tile_window_with_static_distribution return jj == Traits::VectorDimY ? (idx_ys_start[jj] + j) : idx_ys_start[jj]; }, - number{}); + number{}); constexpr index_t d = tile_dstr.get_ys_to_d_descriptor().calculate_offset(idx_ys) / Traits::PackedSize; - vec_value.template get_as()(j / Traits::PackedSize) = + vec_value.template get_as()(j / Traits::PackedSize) = dstr_tensor.get_thread_buffer().template at(); }); // write into bottom tensor - get_bottom_tensor_view() + this->get_bottom_tensor_view() .template set_vectorized_elements_raw( bottom_tensor_thread_coord, 0, vec_value); @@ -729,10 +530,10 @@ struct tile_window_with_static_distribution constexpr auto idx_diff_ys = SFC_Ys::get_forward_step(iAccess); constexpr auto idx_diff_ps_ys = container_concat( - generate_tuple([&](auto) { return number<0>{}; }, number{}), + generate_tuple([&](auto) { return number<0>{}; }, number{}), idx_diff_ys); - move_window_adaptor_and_bottom_tensor_thread_coordinate( + Base::move_window_adaptor_and_bottom_tensor_thread_coordinate( window_adaptor_thread_coord, bottom_tensor_thread_coord, idx_diff_ps_ys); } }); @@ -740,16 +541,18 @@ struct tile_window_with_static_distribution } template - CK_TILE_DEVICE void update(const static_distributed_tensor& dstr_tensor, - number = {}, - bool_constant = {}) const + CK_TILE_DEVICE void + update(const static_distributed_tensor& + dstr_tensor, + number = {}, + bool_constant = {}) const { - using Traits = load_store_traits; + using Traits = typename Base::Traits; using vector_t = typename Traits::vector_t; using SFC_Ys = typename Traits::SFC_Ys; - constexpr auto tile_dstr = TileDstr{}; + constexpr auto tile_dstr = typename Base::TileDstr{}; // loop over thread tensor space [y0, y1, ...] static_for<0, NumCoord, 1>{}([&](auto iCoord) { @@ -772,18 +575,18 @@ struct tile_window_with_static_distribution return jj == Traits::VectorDimY ? (idx_ys_start[jj] + j) : idx_ys_start[jj]; }, - number{}); + number{}); constexpr index_t d = tile_dstr.get_ys_to_d_descriptor().calculate_offset(idx_ys) / Traits::PackedSize; - vec_value.template get_as()(j / Traits::PackedSize) = + vec_value.template get_as()(j / Traits::PackedSize) = dstr_tensor.get_thread_buffer().template at(); }); // write into bottom tensor - get_bottom_tensor_view().template update_vectorized_elements( + this->get_bottom_tensor_view().template update_vectorized_elements( bottom_tensor_thread_coord, 0, vec_value, @@ -795,10 +598,10 @@ struct tile_window_with_static_distribution constexpr auto idx_diff_ys = SFC_Ys::get_forward_step(iAccess); constexpr auto idx_diff_ps_ys = container_concat( - generate_tuple([&](auto) { return number<0>{}; }, number{}), + generate_tuple([&](auto) { return number<0>{}; }, number{}), idx_diff_ys); - move_window_adaptor_and_bottom_tensor_thread_coordinate( + Base::move_window_adaptor_and_bottom_tensor_thread_coordinate( window_adaptor_thread_coord, bottom_tensor_thread_coord, idx_diff_ps_ys); } }); @@ -806,17 +609,19 @@ struct tile_window_with_static_distribution } template - CK_TILE_DEVICE void update_raw(const static_distributed_tensor& dstr_tensor, - number = {}, - bool_constant = {}, - bool_constant = {}) const + CK_TILE_DEVICE void + update_raw(const static_distributed_tensor& + dstr_tensor, + number = {}, + bool_constant = {}, + bool_constant = {}) const { - using Traits = load_store_traits; + using Traits = typename Base::Traits; using vector_t = typename Traits::vector_t; using SFC_Ys = typename Traits::SFC_Ys; - constexpr auto tile_dstr = TileDstr{}; + constexpr auto tile_dstr = typename Base::TileDstr{}; // loop over thread tensor space [y0, y1, ...] static_for<0, NumCoord, 1>{}([&](auto iCoord) { @@ -839,18 +644,18 @@ struct tile_window_with_static_distribution return jj == Traits::VectorDimY ? (idx_ys_start[jj] + j) : idx_ys_start[jj]; }, - number{}); + number{}); constexpr index_t d = tile_dstr.get_ys_to_d_descriptor().calculate_offset(idx_ys) / Traits::PackedSize; - vec_value.template get_as()(j / Traits::PackedSize) = + vec_value.template get_as()(j / Traits::PackedSize) = dstr_tensor.get_thread_buffer().template at(); }); // write into bottom tensor - get_bottom_tensor_view().template update_vectorized_elements_raw( + this->get_bottom_tensor_view().template update_vectorized_elements_raw( bottom_tensor_thread_coord, 0, vec_value, @@ -863,70 +668,44 @@ struct tile_window_with_static_distribution constexpr auto idx_diff_ys = SFC_Ys::get_forward_step(iAccess); constexpr auto idx_diff_ps_ys = container_concat( - generate_tuple([&](auto) { return number<0>{}; }, number{}), + generate_tuple([&](auto) { return number<0>{}; }, number{}), idx_diff_ys); - move_window_adaptor_and_bottom_tensor_thread_coordinate( + Base::move_window_adaptor_and_bottom_tensor_thread_coordinate( window_adaptor_thread_coord, bottom_tensor_thread_coord, idx_diff_ps_ys); } }); }); } - // move thread's botom tensor coordiante - // [x0', x1', ... ] ==> [offset] - // also move window-origin - CK_TILE_DEVICE void move(const BottomTensorIndex& step) + // Custom move behavior + CK_TILE_DEVICE void move_extended(const typename Base::BottomTensorIndex& step) { - window_origin_ += step; - static_for<0, NumCoord, 1>{}([&](auto iCoord) { - move_tensor_coordinate(bottom_tensor_view_.get_tensor_descriptor(), + move_tensor_coordinate(this->bottom_tensor_view_.get_tensor_descriptor(), pre_computed_coords_(iCoord)(I1), step); }); } - CK_TILE_DEVICE void set_window_origin(const BottomTensorIndex& new_window_origin) + CK_TILE_DEVICE void set_window_origin_extended(const typename Base::BottomTensorIndex&) { - window_origin_ = new_window_origin; - -#if 0 // debug - // TODO: this use more register for FA, but less register for GEMM - // need investigation - // only support warp-tile and block-tile - static_assert(NDimP == 1 or NDimP == 2, "wrong!"); - - WindowAdaptorCoord window_adaptor_thread_coord_tmp; - - if constexpr(NDimP == 1) - { - window_adaptor_thread_coord_tmp = make_tensor_adaptor_coordinate( - tile_dstr_.get_ps_ys_to_xs_adaptor(), AdaptorTopIndex{get_lane_id(), 0}); - } - else if constexpr(NDimP == 2) - { - window_adaptor_thread_coord_tmp = - make_tensor_adaptor_coordinate(tile_dstr_.get_ps_ys_to_xs_adaptor(), - AdaptorTopIndex{get_warp_id(), get_lane_id(), 0}); - } -#else // TODO: this use less register for FA, but more register for GEMM // need investigation const auto window_adaptor_thread_coord_tmp = make_tensor_adaptor_coordinate( - tile_dstr_.get_ps_ys_to_xs_adaptor(), - container_concat(detail::get_partition_index(tile_dstr_), array{0})); -#endif + this->tile_dstr_.get_ps_ys_to_xs_adaptor(), + container_concat(detail::get_partition_index(this->tile_dstr_), + array{0})); - BottomTensorIndex bottom_tensor_thread_origin_idx_tmp = - window_origin_ + window_adaptor_thread_coord_tmp.get_bottom_index(); + typename Base::BottomTensorIndex bottom_tensor_thread_origin_idx_tmp = + this->window_origin_ + window_adaptor_thread_coord_tmp.get_bottom_index(); const auto bottom_tensor_thread_coord_tmp = make_tensor_coordinate( - bottom_tensor_view_.get_tensor_descriptor(), bottom_tensor_thread_origin_idx_tmp); + this->bottom_tensor_view_.get_tensor_descriptor(), bottom_tensor_thread_origin_idx_tmp); // pre-compute NumCoord (WindowAdaptorCoord, BottomTensorCoord) bundles to speed up // future load/store() calls (might allocate more registers) - using Traits = load_store_traits; + using Traits = typename Base::Traits; using SFC_Ys = typename Traits::SFC_Ys; static_for<0, NumCoord, 1>{}([&](auto iCoord) { @@ -937,9 +716,10 @@ struct tile_window_with_static_distribution SFC_Ys::get_step_between(number<0>{}, number{}); constexpr auto idx_diff_ps_ys = container_concat( - generate_tuple([&](auto) { return number<0>{}; }, number{}), idx_diff_ys); + generate_tuple([&](auto) { return number<0>{}; }, number{}), + idx_diff_ys); - move_window_adaptor_and_bottom_tensor_thread_coordinate( + Base::move_window_adaptor_and_bottom_tensor_thread_coordinate( window_adaptor_thread_coord, bottom_tensor_thread_coord, idx_diff_ps_ys); pre_computed_coords_(iCoord) = @@ -947,27 +727,11 @@ struct tile_window_with_static_distribution }); } - CK_TILE_HOST_DEVICE void init_raw() { bottom_tensor_view_.init_raw(); } - - // this is the bottom tensor view - // [x0', x1', ...] ==> [offset] - BottomTensorView bottom_tensor_view_; - - // - WindowLengths window_lengths_; - - // origin ([x0', x1', ...]) of window on bottom tensor - BottomTensorIndex window_origin_; - - // Tile tensor distribution, which contains: - // 1. adaptor for window: [p0, p1, ..., y0, y1, ...] ==> [x0, x1, ...] - // 2. thread descriptor for thread tensor in register: [y0, y1, ...] ==> [d] - TileDstr tile_dstr_; - // this contains: // per-thread coordinate for window adaptor // per-thread coordinate for bottom tensor - array, NumCoord> pre_computed_coords_; + array, NumCoord> + pre_computed_coords_; }; // TODO: use strategy @@ -1037,62 +801,26 @@ CK_TILE_DEVICE void move_tile_window( */ template struct tile_window_with_static_lengths + : public tile_window_base, + BottomTensorView_, + WindowLengths_> { - using BottomTensorView = remove_reference_t; - using WindowLengths = remove_cvref_t; - using BottomTensorDesc = typename BottomTensorView::TensorDesc; - using DataType = typename BottomTensorView::DataType; - - static constexpr index_t NDimBottomTensor = BottomTensorDesc::get_num_of_dimension(); - - static_assert(ck_tile::is_known_at_compile_time::value, - "wrong! lengths should be static"); - - using BottomTensorIndex = array; + using Base = + tile_window_base, + BottomTensorView_, + WindowLengths_>; CK_TILE_DEVICE constexpr tile_window_with_static_lengths() = default; CK_TILE_DEVICE constexpr tile_window_with_static_lengths( - const BottomTensorView& bottom_tensor_view, - const WindowLengths& window_lengths, - const BottomTensorIndex& window_origin) - : bottom_tensor_view_{bottom_tensor_view}, - window_lengths_{window_lengths}, - window_origin_{window_origin} + const typename Base::BottomTensorView& bottom_tensor_view, + const typename Base::WindowLengths& window_lengths, + const typename Base::BottomTensorIndex& window_origin) { + this->window_origin_ = window_origin; + this->window_lengths_ = window_lengths; + this->bottom_tensor_view_ = bottom_tensor_view; } - - CK_TILE_DEVICE static constexpr index_t get_num_of_dimension() { return NDimBottomTensor; } - - CK_TILE_DEVICE constexpr auto get_window_lengths() const { return window_lengths_; } - - CK_TILE_DEVICE constexpr auto get_bottom_tensor_view() const { return bottom_tensor_view_; } - - CK_TILE_DEVICE constexpr auto get_window_origin() const { return window_origin_; } - - CK_TILE_DEVICE void set_window_origin(const BottomTensorIndex& new_window_origin) - { - window_origin_ = new_window_origin; - } - - CK_TILE_DEVICE constexpr void - set_bottom_tensor_view_data_ptr(typename BottomTensorView::DataType* data) - { - bottom_tensor_view_.buf_.p_data_ = data; - } - - // move window-origin - CK_TILE_DEVICE void move(const BottomTensorIndex& step) { window_origin_ += step; } - - // this is the bottom tensor view - // [x0', x1', ...] ==> [offset] - BottomTensorView bottom_tensor_view_; - - // - WindowLengths window_lengths_; - - // origin ([x0', x1', ...]) of window on bottom tensor - BottomTensorIndex window_origin_; }; template diff --git a/include/ck_tile/core/tensor/tile_window_base.hpp b/include/ck_tile/core/tensor/tile_window_base.hpp new file mode 100644 index 0000000000..89a928a53c --- /dev/null +++ b/include/ck_tile/core/tensor/tile_window_base.hpp @@ -0,0 +1,256 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved. + +#pragma once + +#include "ck_tile/core/arch/utility.hpp" +#include "ck_tile/core/algorithm/space_filling_curve.hpp" +#include "ck_tile/core/config.hpp" +#include "ck_tile/core/container/array.hpp" +#include "ck_tile/core/container/sequence.hpp" +#include "ck_tile/core/container/tuple.hpp" +#include "ck_tile/core/container/container_helper.hpp" +#include "ck_tile/core/tensor/static_distributed_tensor.hpp" +#include "ck_tile/core/tensor/tensor_adaptor.hpp" +#include "ck_tile/core/tensor/tile_distribution.hpp" +#include "ck_tile/core/utility/functional.hpp" +#include "ck_tile/core/utility/type_traits.hpp" + +namespace ck_tile { + +/** + * @brief This class provides description of tile windowed view on the device memory. + * + * @note This class does not provide any functions to read or modify device memory. + * + * @tparam BottomTensorView_ Class describing & holding device tensor memory. + * @tparam WindowLengths_ Spatial sizes of windowed view on tensor. + */ +template +struct tile_window_base +{ + + using BottomTensorView = remove_reference_t; + using WindowLengths = remove_cvref_t; + using BottomTensorDesc = typename BottomTensorView::TensorDesc; + using DataType = remove_cvref_t; + + static constexpr index_t NDimBottomTensor = BottomTensorDesc::get_num_of_dimension(); + + static_assert(ck_tile::is_known_at_compile_time::value, + "wrong! lengths should be static"); + + using BottomTensorIndex = array; + + CK_TILE_DEVICE constexpr auto get_window_origin() const { return window_origin_; } + CK_TILE_DEVICE constexpr auto get_window_lengths() const { return window_lengths_; } + CK_TILE_DEVICE constexpr auto get_bottom_tensor_view() const { return bottom_tensor_view_; } + CK_TILE_DEVICE static constexpr index_t get_num_of_dimension() { return NDimBottomTensor; } + + CK_TILE_DEVICE void set_window_origin(const BottomTensorIndex& new_window_origin) + { + window_origin_ = new_window_origin; + + // Delegate to child if it implements extra logic + static_cast(this)->set_window_origin_extended(new_window_origin); + } + // Default no-op; can be overridden in child + CK_TILE_DEVICE void set_window_origin_extended(const BottomTensorIndex&) {} + + CK_TILE_DEVICE constexpr void + set_bottom_tensor_view_data_ptr(typename BottomTensorView::DataType* data) + { + bottom_tensor_view_.buf_.p_data_ = data; + } + + // move window-origin + CK_TILE_DEVICE void move(const BottomTensorIndex& step) + { + window_origin_ += step; + + // Delegate to child if it implements extra movement logic + static_cast(this)->move_extended(step); + } + + // Default no-op; can be overridden in child + CK_TILE_DEVICE void move_extended(const BottomTensorIndex&) {} + + // origin ([x0', x1', ...]) of window on bottom tensor + BottomTensorIndex window_origin_; + + WindowLengths window_lengths_; + + // this is the bottom tensor view + // [x0', x1', ...] ==> [offset] + BottomTensorView bottom_tensor_view_; +}; + +template +struct tile_window_with_tile_dstr_base + : public tile_window_base +{ + using TileDstr = remove_cvref_t; + using TileWindowBase = tile_window_base; + + using WindowAdaptor = typename TileDstr::PsYs2XsAdaptor; + + static constexpr index_t NDimWindowAdaptorTop = WindowAdaptor::get_num_of_top_dimension(); + + static constexpr index_t NDimP = TileDstr::get_num_of_dimension_p(); + static constexpr index_t NDimY = TileDstr::get_num_of_dimension_y(); + + using AdaptorTopIndex = array; + // using BottomTensorIndex = array; + + using WindowAdaptorCoord = + decltype(make_tensor_adaptor_coordinate(WindowAdaptor{}, AdaptorTopIndex{})); + + using BottomTensorCoord = decltype(make_tensor_coordinate( + typename TileWindowBase::BottomTensorDesc{}, typename TileWindowBase::BottomTensorIndex{})); + + static_assert(TileDstr::is_static(), "wrong!"); + static_assert(TileWindowBase::NDimBottomTensor == WindowAdaptor::get_num_of_bottom_dimension(), + "wrong! inconsistent # of diemsnions"); + + CK_TILE_DEVICE constexpr auto get_tile_distribution() const { return tile_dstr_; } + CK_TILE_HOST_DEVICE void init_raw() { this->bottom_tensor_view_.init_raw(); } + + CK_TILE_DEVICE static constexpr bool has_static_tile_distribution() + { + return TileDstr::is_static(); + } + + // move thread's window adaptor coordinate and bottom tensor coordinate + // [p0, p1, ..., y0, y1, ...] ==> [x0, x1, ...] ==> [x0', x1', ...] ==> [offset] + template + CK_TILE_DEVICE void move_window_adaptor_and_bottom_tensor_thread_coordinate( + WindowAdaptorCoord& window_adaptor_thread_coord, + BottomTensorCoord& bottom_tensor_thread_coord, + const ATopIndex& idx_diff_adaptor_top) const + { + array idx_diff_adaptor_bottom; + + move_tensor_adaptor_coordinate(tile_dstr_.get_ps_ys_to_xs_adaptor(), + window_adaptor_thread_coord, + idx_diff_adaptor_top, + idx_diff_adaptor_bottom); + + move_tensor_coordinate(this->bottom_tensor_view_.get_tensor_descriptor(), + bottom_tensor_thread_coord, + idx_diff_adaptor_bottom); + } + + struct Traits + { + public: + static constexpr index_t PackedSize = + ck_tile::numeric_traits>::PackedSize; + + static constexpr auto get_vector_dim_y_scalar_per_vector() + { + const auto [ys_vector_lengths, ys_vector_strides] = + tile_window_with_tile_dstr_base::get_window_adaptor_ys_safe_vector_length_strides(); + + index_t VectorDimY_ = 0; + index_t ScalarPerVector_ = 1; + + for(index_t i = 0; i < NDimY; ++i) + { + if(ys_vector_strides[i] == 1 && ys_vector_lengths[i] > ScalarPerVector_) + { + ScalarPerVector_ = ys_vector_lengths[i]; + VectorDimY_ = i; + } + } + + return make_tuple(VectorDimY_, ScalarPerVector_); + } + + static constexpr index_t VectorDimY = get_vector_dim_y_scalar_per_vector().template at<0>(); + static constexpr index_t ScalarPerVector = + get_vector_dim_y_scalar_per_vector().template at<1>(); + using vector_t = + thread_buffer; + + static constexpr auto scalars_per_access_ = [] { + constexpr auto scalars_per_access_arr = generate_array( + [&](auto i) { return (i == VectorDimY) ? ScalarPerVector : 1; }, number{}); + + /// TODO: add non-automatic storage argument support to macro TO_SEQUENCE() + constexpr auto NDimY_ = NDimY; + + return TO_SEQUENCE(scalars_per_access_arr, NDimY_); + }(); + + static constexpr auto get_space_filling_curve() + { + constexpr auto thread_tensor_lengths_ys = + to_sequence(TileDstr{}.get_ys_to_d_descriptor().get_lengths()); + + // FIXME: need logic to judge dim access order + using DimAccessOrder = typename arithmetic_sequence_gen<0, NDimY, 1>::type; + + return space_filling_curve{}; + } + + using SFC_Ys = decltype(get_space_filling_curve()); + + static constexpr index_t NumAccess = SFC_Ys::get_num_of_access(); + + static_assert(0 < NumAccess, "Wrong! NumAccess should be larger than 0"); + }; + + // return vector dimension among [y0, y1, ...] + CK_TILE_DEVICE static constexpr auto get_window_adaptor_ys_safe_vector_length_strides() + { + // bottom tensor top dimension vector lengths and strides + const auto [bottom_tensor_top_dim_vector_lengths, bottom_tensor_top_dim_vector_strides] = + TileWindowBase::BottomTensorDesc::get_top_dimension_safe_vector_length_strides(); + + // window vector lengths/strides + const auto window_adaptor_bottom_dim_vector_lengths = bottom_tensor_top_dim_vector_lengths; + const auto window_adaptor_bottom_dim_vector_strides = bottom_tensor_top_dim_vector_strides; + + // window adaptor [p0, p1, ..., y0, y1, ...] + array window_adaptor_vector_lengths{ + -1}; + array window_adaptor_vector_strides{ + -1}; + + constexpr auto window_adaptor_bottom_dims = + WindowAdaptor::get_bottom_dimension_hidden_ids(); + + set_container_subset(window_adaptor_vector_lengths, + window_adaptor_bottom_dims, + window_adaptor_bottom_dim_vector_lengths); + set_container_subset(window_adaptor_vector_strides, + window_adaptor_bottom_dims, + window_adaptor_bottom_dim_vector_strides); + + const auto [window_adaptor_ps_ys_vector_lengths, window_adaptor_ps_ys_vector_strides] = + WindowAdaptor{}.get_top_dimension_safe_vector_length_strides( + window_adaptor_vector_lengths, window_adaptor_vector_strides); + + // [y0, y1, ...] + constexpr auto y_dims = typename arithmetic_sequence_gen::type{}; + + return make_tuple(get_container_subset(window_adaptor_ps_ys_vector_lengths, y_dims), + get_container_subset(window_adaptor_ps_ys_vector_strides, y_dims)); + } + + CK_TILE_DEVICE constexpr auto get_num_of_access() const { return Traits::NumAccess; } + // Tile tensor distribution, which contains: + // 1. adaptor for window: [p0, p1, ..., y0, y1, ...] ==> [x0, x1, ...] + // 2. thread descriptor for thread tensor in register: [y0, y1, ...] ==> [d] + TileDstr tile_dstr_; +}; + +} // namespace ck_tile diff --git a/include/ck_tile/core/tensor/tile_window_linear.hpp b/include/ck_tile/core/tensor/tile_window_linear.hpp index 5ecaf5ca17..f11610d658 100644 --- a/include/ck_tile/core/tensor/tile_window_linear.hpp +++ b/include/ck_tile/core/tensor/tile_window_linear.hpp @@ -13,6 +13,7 @@ #include "ck_tile/core/tensor/static_distributed_tensor.hpp" #include "ck_tile/core/tensor/tensor_adaptor.hpp" #include "ck_tile/core/tensor/tile_distribution.hpp" +#include "ck_tile/core/tensor/tile_window_base.hpp" #include "ck_tile/core/utility/functional.hpp" #include "ck_tile/core/utility/type_traits.hpp" @@ -37,171 +38,48 @@ namespace ck_tile { // TODO: if using this struct, better use load_raw()/store_raw(), can control // the the immediate offset on the fly // space-filing-curve is non-snaked here! -// +// This struct inherits from tile_window_with_tile_dstr_base, which is an intermediary base class +// with the ultimate parent class being tile_window_base. template struct tile_window_linear + : public tile_window_with_tile_dstr_base, + BottomTensorView_, + WindowLengths_, + StaticTileDistribution_> { + using Base = tile_window_with_tile_dstr_base, + BottomTensorView_, + WindowLengths_, + StaticTileDistribution_>; - using BottomTensorView = remove_reference_t; - using WindowLengths = remove_cvref_t; - using TileDstr = remove_cvref_t; - - using WindowAdaptor = typename TileDstr::PsYs2XsAdaptor; - using BottomTensorDesc = typename BottomTensorView::TensorDesc; - - using DataType = remove_cvref_t; using LinearBottomDims = remove_cvref_t; - static_assert(LinearBottomDims::size() == BottomTensorView::get_num_of_dimension()); - - static constexpr index_t NDimWindowAdaptorTop = WindowAdaptor::get_num_of_top_dimension(); - static constexpr index_t NDimBottomTensor = BottomTensorDesc::get_num_of_dimension(); - - static constexpr index_t NDimP = TileDstr::get_num_of_dimension_p(); - static constexpr index_t NDimY = TileDstr::get_num_of_dimension_y(); + static_assert(LinearBottomDims::size() == Base::BottomTensorView::get_num_of_dimension()); static constexpr auto I0 = number<0>{}; static constexpr auto I1 = number<1>{}; - // TODO: check WindowLengths and StaticTileDistribution are consistent - - static_assert(ck_tile::is_known_at_compile_time::value, - "wrong! lengths should be static"); - static_assert(TileDstr::is_static(), "wrong!"); - - static_assert(NDimBottomTensor == WindowAdaptor::get_num_of_bottom_dimension(), - "wrong! inconsistent # of diemsnions"); - - using AdaptorTopIndex = array; - using BottomTensorIndex = array; - - using WindowAdaptorCoord = - decltype(make_tensor_adaptor_coordinate(WindowAdaptor{}, AdaptorTopIndex{})); - - using BottomTensorCoord = - decltype(make_tensor_coordinate(BottomTensorDesc{}, BottomTensorIndex{})); - struct traits { - private: - // return vector dimension among [y0, y1, ...] - CK_TILE_DEVICE static constexpr auto get_window_adaptor_ys_safe_vector_length_strides() - { - // bottom tensor top dimension vector lengths and strides - const auto [bottom_tensor_top_dim_vector_lengths, - bottom_tensor_top_dim_vector_strides] = - BottomTensorDesc::get_top_dimension_safe_vector_length_strides(); - - // window vector lengths/strides - const auto window_adaptor_bottom_dim_vector_lengths = - bottom_tensor_top_dim_vector_lengths; - const auto window_adaptor_bottom_dim_vector_strides = - bottom_tensor_top_dim_vector_strides; - - // window adaptor [p0, p1, ..., y0, y1, ...] - array - window_adaptor_vector_lengths{-1}; - array - window_adaptor_vector_strides{-1}; - - constexpr auto window_adaptor_bottom_dims = - WindowAdaptor::get_bottom_dimension_hidden_ids(); - - set_container_subset(window_adaptor_vector_lengths, - window_adaptor_bottom_dims, - window_adaptor_bottom_dim_vector_lengths); - set_container_subset(window_adaptor_vector_strides, - window_adaptor_bottom_dims, - window_adaptor_bottom_dim_vector_strides); - - const auto [window_adaptor_ps_ys_vector_lengths, window_adaptor_ps_ys_vector_strides] = - WindowAdaptor{}.get_top_dimension_safe_vector_length_strides( - window_adaptor_vector_lengths, window_adaptor_vector_strides); - - // [y0, y1, ...] - constexpr auto y_dims = - typename arithmetic_sequence_gen::type{}; - - return make_tuple(get_container_subset(window_adaptor_ps_ys_vector_lengths, y_dims), - get_container_subset(window_adaptor_ps_ys_vector_strides, y_dims)); - } - - static constexpr auto get_vector_dim_y_scalar_per_vector() - { - const auto [ys_vector_lengths, ys_vector_strides] = - get_window_adaptor_ys_safe_vector_length_strides(); - - index_t VectorDimY_ = 0; - index_t ScalarPerVector_ = 1; - - for(index_t i = 0; i < NDimY; ++i) - { - if(ys_vector_strides[i] == 1 && ys_vector_lengths[i] > ScalarPerVector_) - { - ScalarPerVector_ = ys_vector_lengths[i]; - VectorDimY_ = i; - } - } - - return make_tuple(VectorDimY_, ScalarPerVector_); - } - - public: - static constexpr index_t PackedSize = - ck_tile::numeric_traits>::PackedSize; - static constexpr index_t VectorDimY = get_vector_dim_y_scalar_per_vector().template at<0>(); - static constexpr index_t ScalarPerVector = - get_vector_dim_y_scalar_per_vector().template at<1>(); - - using vector_t = thread_buffer; - - private: - static constexpr auto scalars_per_access_ = [] { - constexpr auto scalars_per_access_arr = generate_array( - [&](auto i) { return (i == VectorDimY) ? ScalarPerVector : 1; }, number{}); - - /// TODO: add non-automatic storage argument support to macro TO_SEQUENCE() - constexpr auto NDimY_ = NDimY; - - return TO_SEQUENCE(scalars_per_access_arr, NDimY_); - }(); - - static constexpr auto get_space_filling_curve() - { - constexpr auto thread_tensor_lengths_ys = - to_sequence(TileDstr{}.get_ys_to_d_descriptor().get_lengths()); - - // FIXME: need logic to judge dim access order - using DimAccessOrder = typename arithmetic_sequence_gen<0, NDimY, 1>::type; - - return space_filling_curve{}; - } - - public: - using SFC_Ys = decltype(get_space_filling_curve()); - - static constexpr index_t NumAccess = SFC_Ys::get_num_of_access(); - - static_assert(0 < NumAccess, "Wrong! NumAccess should be larger than 0"); - private: static constexpr auto get_num_non_linear_access() { - constexpr auto sfc_access_lens = SFC_Ys::access_lengths; - using ys_to_rhs_major = - typename decltype(TileDstr{}.get_static_tile_distribution_encoding())::Ys2RHsMajor; + constexpr auto sfc_access_lens = Base::Traits::SFC_Ys::access_lengths; + using ys_to_rhs_major = typename decltype( + typename Base::TileDstr{}.get_static_tile_distribution_encoding())::Ys2RHsMajor; constexpr auto non_linear = [&]() { index_t cnt = 1; - static_for<0, NDimY, 1>{}([&](auto i_dim_y) { + static_for<0, Base::NDimY, 1>{}([&](auto i_dim_y) { constexpr auto rhs_major = ys_to_rhs_major{}[i_dim_y]; constexpr auto target_h_dim = number{}; // no r dim here! if constexpr(LinearBottomDims{}[target_h_dim] == 0) @@ -230,20 +108,20 @@ struct tile_window_linear // -> prefixsum : seqneuce<0, 2, 4, 6, 8> static constexpr auto get_non_linear_access_map() { - constexpr auto sfc_access_lens = SFC_Ys::access_lengths; - using ys_to_rhs_major = - typename decltype(TileDstr{}.get_static_tile_distribution_encoding())::Ys2RHsMajor; + constexpr auto sfc_access_lens = Base::Traits::SFC_Ys::access_lengths; + using ys_to_rhs_major = typename decltype( + typename Base::TileDstr{}.get_static_tile_distribution_encoding())::Ys2RHsMajor; constexpr auto non_linear_map = [&]() { - array m_{0}; + array m_{0}; index_t cumulative_len_ = 1; index_t cumulative_non_linear_len_ = 1; - static_for<0, NDimY, 1>{}([&](auto i_y) { - constexpr auto i_dim_y = number{}; // from right to left + static_for<0, Base::NDimY, 1>{}([&](auto i_y) { + constexpr auto i_dim_y = number{}; // from right to left constexpr auto rhs_major = ys_to_rhs_major{}[i_dim_y]; constexpr auto target_h_dim = number{}; // no r dim here! constexpr auto is_linear_dim = LinearBottomDims{}[target_h_dim]; - array current_m_{0}; + array current_m_{0}; constexpr auto current_len_ = sfc_access_lens[i_dim_y]; // copy cumulative length as current pattern @@ -266,13 +144,12 @@ struct tile_window_linear return m_; }(); - return TO_SEQUENCE(non_linear_map, NumAccess); + return TO_SEQUENCE(non_linear_map, Base::Traits::NumAccess); } static constexpr auto get_non_linear_access_histogram() { constexpr auto m_ = get_non_linear_access_map(); - // m_.foo(); constexpr auto r_ = typename arithmetic_sequence_gen<0, get_num_non_linear_access() + 1, 1>::type{}; @@ -296,7 +173,7 @@ struct tile_window_linear using AccessPrefixSum_NonLinear = decltype(get_non_linear_access_histogram_prefix_sum()); }; - static constexpr index_t NumAccess = traits::NumAccess; + static constexpr index_t NumAccess = Base::Traits::NumAccess; static constexpr index_t NumAccess_NonLinear = traits::NumAccess_NonLinear; using AccessMap_NonLinear = typename traits::AccessMap_NonLinear; using AccessHistogram_NonLinear = typename traits::AccessHistogram_NonLinear; @@ -304,30 +181,31 @@ struct tile_window_linear CK_TILE_DEVICE constexpr tile_window_linear() = default; - CK_TILE_DEVICE constexpr tile_window_linear(const BottomTensorView& bottom_tensor_view, - const WindowLengths& window_lengths, - const BottomTensorIndex& window_origin, - const TileDstr& tile_distribution) - : bottom_tensor_view_{bottom_tensor_view}, - window_lengths_{window_lengths}, - window_origin_{window_origin}, - tile_dstr_{tile_distribution}, - cached_coords_{}, - cached_flags_{} + CK_TILE_DEVICE constexpr tile_window_linear( + const typename Base::BottomTensorView& bottom_tensor_view, + const typename Base::WindowLengths& window_lengths, + const typename Base::BottomTensorIndex& window_origin, + const typename Base::TileDstr& tile_distribution) + : cached_coords_{}, cached_flags_{} { + this->bottom_tensor_view_ = bottom_tensor_view; + this->window_lengths_ = window_lengths; + this->window_origin_ = window_origin; + this->tile_dstr_ = tile_distribution; auto window_adaptor_thread_coord_tmp = make_tensor_adaptor_coordinate( tile_distribution.get_ps_ys_to_xs_adaptor(), - container_concat(make_tuple(get_warp_id(), get_lane_id()), - generate_tuple([&](auto) { return number<0>{}; }, number{}))); + container_concat( + make_tuple(get_warp_id(), get_lane_id()), + generate_tuple([&](auto) { return number<0>{}; }, number{}))); - BottomTensorIndex bottom_tensor_thread_origin_idx_tmp = + typename Base::BottomTensorIndex bottom_tensor_thread_origin_idx_tmp = window_origin + window_adaptor_thread_coord_tmp.get_bottom_index(); auto bottom_tensor_thread_coord_tmp = make_tensor_coordinate( - bottom_tensor_view_.get_tensor_descriptor(), bottom_tensor_thread_origin_idx_tmp); + this->bottom_tensor_view_.get_tensor_descriptor(), bottom_tensor_thread_origin_idx_tmp); // future load/store() calls (might allocate more registers) - using SFC_Ys = typename traits::SFC_Ys; + using SFC_Ys = typename Base::Traits::SFC_Ys; static_for<0, NumAccess, 1>{}([&](auto i_access) { constexpr auto non_linear_id = number{}; @@ -343,16 +221,16 @@ struct tile_window_linear // cached flag is independent from non-linear-coord // but need be updated in move_tile, with proper dims cached_flags_(i_access) = coordinate_has_valid_offset_assuming_top_index_is_valid( - bottom_tensor_view_.get_tensor_descriptor(), bottom_tensor_thread_coord_tmp); + this->bottom_tensor_view_.get_tensor_descriptor(), bottom_tensor_thread_coord_tmp); if constexpr(i_access != (NumAccess - 1)) { constexpr auto idx_diff_ys = SFC_Ys::get_forward_step(i_access); // tuple of number constexpr auto idx_diff_ps_ys = container_concat( - generate_tuple([&](auto) { return number<0>{}; }, number{}), + generate_tuple([&](auto) { return number<0>{}; }, number{}), idx_diff_ys); - move_window_adaptor_and_bottom_tensor_thread_coordinate( + Base::move_window_adaptor_and_bottom_tensor_thread_coordinate( window_adaptor_thread_coord_tmp, bottom_tensor_thread_coord_tmp, idx_diff_ps_ys); @@ -360,54 +238,13 @@ struct tile_window_linear }); } - CK_TILE_DEVICE static constexpr index_t get_num_of_dimension() { return NDimBottomTensor; } - - CK_TILE_DEVICE static constexpr bool has_static_tile_distribution() - { - return TileDstr::is_static(); - } - - CK_TILE_DEVICE constexpr auto get_window_lengths() const { return window_lengths_; } - - CK_TILE_DEVICE constexpr auto get_tile_distribution() const { return tile_dstr_; } - - CK_TILE_DEVICE constexpr auto get_bottom_tensor_view() const { return bottom_tensor_view_; } - - CK_TILE_DEVICE constexpr auto get_window_origin() const { return window_origin_; } - - CK_TILE_DEVICE constexpr void - set_bottom_tensor_view_data_ptr(typename BottomTensorView::DataType* data) - { - bottom_tensor_view_.buf_.p_data_ = data; - } - - // move thread's window adaptor coordinate and bottom tensor coordinate - // [p0, p1, ..., y0, y1, ...] ==> [x0, x1, ...] ==> [x0', x1', ...] ==> [offset] - template - CK_TILE_DEVICE void move_window_adaptor_and_bottom_tensor_thread_coordinate( - WindowAdaptorCoord& window_adaptor_thread_coord, - BottomTensorCoord& bottom_tensor_thread_coord, - const ATopIndex& idx_diff_adaptor_top) const - { - array idx_diff_adaptor_bottom; - - move_tensor_adaptor_coordinate(tile_dstr_.get_ps_ys_to_xs_adaptor(), - window_adaptor_thread_coord, - idx_diff_adaptor_top, - idx_diff_adaptor_bottom); - - move_tensor_coordinate(bottom_tensor_view_.get_tensor_descriptor(), - bottom_tensor_thread_coord, - idx_diff_adaptor_bottom); - } - template CK_TILE_DEVICE static constexpr auto get_bottom_linear_coordinate(number) { - using SFC_Ys = typename traits::SFC_Ys; + using SFC_Ys = typename Base::Traits::SFC_Ys; constexpr auto idx_ys = SFC_Ys::get_index(number{}); - using ys_to_rhs_major = - typename decltype(TileDstr{}.get_static_tile_distribution_encoding())::Ys2RHsMajor; + using ys_to_rhs_major = typename decltype( + typename Base::TileDstr{}.get_static_tile_distribution_encoding())::Ys2RHsMajor; constexpr auto modified_idx_ys = generate_tuple( [&](auto i_dim_y) { @@ -422,9 +259,9 @@ struct tile_window_linear return number{}; } }, - number{}); + number{}); - constexpr auto adaptor_ = TileDstr{}.get_ps_ys_to_xs_adaptor(); + constexpr auto adaptor_ = typename Base::TileDstr{}.get_ps_ys_to_xs_adaptor(); constexpr auto idx_ = container_concat(make_tuple(number<0>{}, number<0>{}), modified_idx_ys); @@ -441,8 +278,8 @@ struct tile_window_linear { // this case usually is a LDS window, everything is known at compile tile. // we directly use BottomTensorView transform to compute the offset, in case padding - auto bottom_tensor_coord = - make_tensor_coordinate(BottomTensorView{}.get_tensor_descriptor(), linear_coord); + auto bottom_tensor_coord = make_tensor_coordinate( + typename Base::BottomTensorView{}.get_tensor_descriptor(), linear_coord); return bottom_tensor_coord.get_offset(); } else @@ -453,7 +290,7 @@ struct tile_window_linear // since that would introduce runtime length (so can't use linear offset) constexpr index_t linear_offset = [&]() { constexpr auto x_idx_ = linear_coord; - constexpr auto x_len_ = TileDstr{}.get_lengths(); + constexpr auto x_len_ = typename Base::TileDstr{}.get_lengths(); static_assert(x_idx_.size() == x_len_.size()); constexpr index_t x_dims_ = x_idx_.size(); index_t cu_stride_ = 1; @@ -469,17 +306,16 @@ struct tile_window_linear } } - CK_TILE_DEVICE constexpr auto get_num_of_access() const { return traits::NumAccess; } - template CK_TILE_DEVICE auto load(number = {}, bool_constant = {}) const { - using vector_t = typename traits::vector_t; - using SFC_Ys = typename traits::SFC_Ys; + using vector_t = typename Base::Traits::vector_t; + using SFC_Ys = typename Base::Traits::SFC_Ys; - constexpr auto tile_dstr = TileDstr{}; + constexpr auto tile_dstr = typename Base::TileDstr{}; - auto dst_tensor = make_static_distributed_tensor(tile_dstr); + auto dst_tensor = + make_static_distributed_tensor(tile_dstr); auto issue = [&](auto i_access_) { constexpr auto IAccess = number{}; @@ -492,35 +328,29 @@ struct tile_window_linear // read from bottom tensor const vector_t vec_value = - get_bottom_tensor_view().template get_vectorized_elements( + this->get_bottom_tensor_view().template get_vectorized_elements( bottom_tensor_thread_coord, linear_offset, bottom_tensor_flag, bool_constant{}); -#if 1 + // data index [y0, y1, ...] constexpr auto idx_diff_ys = SFC_Ys::get_index(IAccess); // write into distributed tensor - static_for<0, traits::ScalarPerVector, traits::PackedSize>{}([&](auto j) { + static_for<0, Base::Traits::ScalarPerVector, Base::Traits::PackedSize>{}([&](auto j) { constexpr auto idx_ys = generate_tuple( [&](auto jj) { - return jj == traits::VectorDimY ? (idx_diff_ys[jj] + j) : idx_diff_ys[jj]; + return jj == Base::Traits::VectorDimY ? (idx_diff_ys[jj] + j) + : idx_diff_ys[jj]; }, - number{}); + number{}); constexpr index_t d = tile_dstr.get_ys_to_d_descriptor().calculate_offset(idx_ys) / - traits::PackedSize; + Base::Traits::PackedSize; - dst_tensor.get_thread_buffer().template at() = - vec_value.template get_as()[j / traits::PackedSize]; + dst_tensor.get_thread_buffer().template at() = vec_value.template get_as< + typename Base::DataTypeDataType>()[j / Base::Traits::PackedSize]; }); -#else - constexpr index_t d = tile_dstr.get_ys_to_d_descriptor().calculate_offset(idx_ys_start); - static_assert(d % traits::ScalarPerVector == 0); - - dst_tensor.get_thread_buffer().template get_as()( - number{}) = bit_cast(vec_value); -#endif }; WINDOW_DISPATCH_ISSUE(); @@ -533,10 +363,10 @@ struct tile_window_linear number = {}, bool_constant = {}) const { - using vector_t = typename traits::vector_t; - using SFC_Ys = typename traits::SFC_Ys; + using vector_t = typename Base::Traits::vector_t; + using SFC_Ys = typename Base::Traits::SFC_Ys; - constexpr auto tile_dstr = TileDstr{}; + constexpr auto tile_dstr = typename Base::TileDstr{}; // auto dst_tensor = make_static_distributed_tensor(tile_dstr); @@ -551,35 +381,28 @@ struct tile_window_linear // read from bottom tensor const vector_t vec_value = - get_bottom_tensor_view().template get_vectorized_elements( + this->get_bottom_tensor_view().template get_vectorized_elements( bottom_tensor_thread_coord, linear_offset, bottom_tensor_flag, bool_constant{}); -#if 1 // data index [y0, y1, ...] constexpr auto idx_diff_ys = SFC_Ys::get_index(IAccess); // write into distributed tensor - static_for<0, traits::ScalarPerVector, traits::PackedSize>{}([&](auto j) { + static_for<0, Base::Traits::ScalarPerVector, Base::Traits::PackedSize>{}([&](auto j) { constexpr auto idx_ys = generate_tuple( [&](auto jj) { - return jj == traits::VectorDimY ? (idx_diff_ys[jj] + j) : idx_diff_ys[jj]; + return jj == Base::Traits::VectorDimY ? (idx_diff_ys[jj] + j) + : idx_diff_ys[jj]; }, - number{}); + number{}); constexpr index_t d = tile_dstr.get_ys_to_d_descriptor().calculate_offset(idx_ys) / - traits::PackedSize; + Base::Traits::PackedSize; - dst_tensor.get_thread_buffer().template at() = - vec_value.template get_as()[j / traits::PackedSize]; + dst_tensor.get_thread_buffer().template at() = vec_value.template get_as< + typename Base::DataTypeDataType>()[j / Base::Traits::PackedSize]; }); -#else - constexpr index_t d = tile_dstr.get_ys_to_d_descriptor().calculate_offset(idx_ys_start); - static_assert(d % traits::ScalarPerVector == 0); - - dst_tensor.get_thread_buffer().template get_as()( - number{}) = bit_cast(vec_value); -#endif }; WINDOW_DISPATCH_ISSUE(); @@ -596,15 +419,17 @@ struct tile_window_linear bool_constant = {}, bool_constant = {}) const { - using vector_t = typename traits::vector_t; - using SFC_Ys = typename traits::SFC_Ys; + using vector_t = typename Base::Traits::vector_t; + using SFC_Ys = typename Base::Traits::SFC_Ys; static constexpr index_t YElementSize = - TileDstr{}.get_ys_to_d_descriptor().get_element_space_size(); - static_assert(YElementSize % (traits::PackedSize * traits::ScalarPerVector) == 0); + typename Base::TileDstr{}.get_ys_to_d_descriptor().get_element_space_size(); + static_assert(YElementSize % (Base::Traits::PackedSize * Base::Traits::ScalarPerVector) == + 0); using vectorized_tbuf = - array; + array; - constexpr auto tile_dstr = TileDstr{}; + constexpr auto tile_dstr = typename Base::TileDstr{}; auto& dst_vec_tbuf = reinterpret_cast(dst_tensor.get_thread_buffer()); @@ -612,7 +437,7 @@ struct tile_window_linear constexpr auto IAccess = number{}; constexpr auto pre_nop_ = [&]() { if constexpr(pre_nop && i_access_ == 0 && - BottomTensorView::buffer_view::get_address_space() == + Base::BottomTensorView::buffer_view::get_address_space() == address_space_enum::global) return bool_constant{}; else @@ -628,11 +453,11 @@ struct tile_window_linear constexpr auto idx_ys_start = SFC_Ys::get_index(IAccess); constexpr index_t d = tile_dstr.get_ys_to_d_descriptor().calculate_offset(idx_ys_start) / - traits::PackedSize; - static_assert(d % traits::ScalarPerVector == 0); + Base::Traits::PackedSize; + static_assert(d % Base::Traits::ScalarPerVector == 0); - get_bottom_tensor_view().template get_vectorized_elements_raw( - dst_vec_tbuf.template at(), + this->get_bottom_tensor_view().template get_vectorized_elements_raw( + dst_vec_tbuf.template at(), bottom_tensor_thread_coord, linear_offset /**/, bottom_tensor_flag, @@ -663,7 +488,7 @@ struct tile_window_linear // currently we only support everything is non linear dim // actually it's not performant if we have linear dim(e.g. fast changing) static_assert(NumAccess_NonLinear == NumAccess); - static_assert(BottomTensorView::buffer_view::get_address_space() == + static_assert(Base::BottomTensorView::buffer_view::get_address_space() == address_space_enum::global); // issues * warps * lanes @@ -689,7 +514,7 @@ struct tile_window_linear const index_t m0_init_value = size_per_buf + size_per_wave * get_warp_id(); m0_set_with_memory(m0_init_value); // This should be wave independent - using vector_t = typename traits::vector_t; + using vector_t = typename Base::Traits::vector_t; LdsDataType* smem = lds_tile.get_bottom_tensor_view().get_buffer_view().p_data_; @@ -708,7 +533,7 @@ struct tile_window_linear auto bottom_tensor_flag = cached_flags_[IAccess]; // get this flag anyway // read from bottom tensor - get_bottom_tensor_view().template async_get_vectorized_elements_raw( + this->get_bottom_tensor_view().template async_get_vectorized_elements_raw( smem, bottom_tensor_thread_coord, 0, bottom_tensor_flag, pre_nop_); // move thread coordinate @@ -732,7 +557,7 @@ struct tile_window_linear // currently we only support everything is non linear dim // actually it's not performant if we have linear dim(e.g. fast changing) static_assert(NumAccess_NonLinear == NumAccess); - static_assert(BottomTensorView::buffer_view::get_address_space() == + static_assert(Base::BottomTensorView::buffer_view::get_address_space() == address_space_enum::global); // issues * warps * lanes @@ -757,7 +582,7 @@ struct tile_window_linear const index_t m0_init_value = size_per_buf + size_per_wave * get_warp_id(); - using vector_t = typename traits::vector_t; + using vector_t = typename Base::Traits::vector_t; // TODO: we force CK_TILE_LDS_ADDR CK_TILE_LDS_ADDR LdsDataType* smem = @@ -771,7 +596,7 @@ struct tile_window_linear auto bottom_tensor_flag = cached_flags_[IAccess]; // read from bottom tensor - get_bottom_tensor_view().template async_get_vectorized_elements( + this->get_bottom_tensor_view().template async_get_vectorized_elements( smem, bottom_tensor_thread_coord, 0, @@ -789,15 +614,16 @@ struct tile_window_linear } template - CK_TILE_DEVICE void store(const static_distributed_tensor& dstr_tensor, + CK_TILE_DEVICE void store(const static_distributed_tensor& dstr_tensor, number = {}, bool_constant = {}) const { - using vector_t = typename traits::vector_t; - using SFC_Ys = typename traits::SFC_Ys; + using vector_t = typename Base::Traits::vector_t; + using SFC_Ys = typename Base::Traits::SFC_Ys; - constexpr auto tile_dstr = TileDstr{}; + constexpr auto tile_dstr = typename Base::TileDstr{}; // loop over thread tensor space [y0, y1, ...] auto issue = [&](auto i_access_) { @@ -812,22 +638,23 @@ struct tile_window_linear // read from distributed tensor vector_t vec_value; - static_for<0, traits::ScalarPerVector, traits::PackedSize>{}([&](auto j) { + static_for<0, Base::Traits::ScalarPerVector, Base::Traits::PackedSize>{}([&](auto j) { constexpr auto idx_ys = generate_tuple( [&](auto jj) { - return jj == traits::VectorDimY ? (idx_ys_start[jj] + j) : idx_ys_start[jj]; + return jj == Base::Traits::VectorDimY ? (idx_ys_start[jj] + j) + : idx_ys_start[jj]; }, - number{}); + number{}); constexpr index_t d = tile_dstr.get_ys_to_d_descriptor().calculate_offset(idx_ys) / - traits::PackedSize; + Base::Traits::PackedSize; - vec_value.template get_as()(j / traits::PackedSize) = + vec_value.template get_as()(j / Base::Traits::PackedSize) = dstr_tensor.get_thread_buffer().template at(); }); // write into bottom tensor - get_bottom_tensor_view().template set_vectorized_elements( + this->get_bottom_tensor_view().template set_vectorized_elements( bottom_tensor_thread_coord, linear_offset, bottom_tensor_flag, @@ -839,13 +666,15 @@ struct tile_window_linear } template - CK_TILE_DEVICE void store_raw(const static_distributed_tensor& dstr_tensor, - number = {}) const + CK_TILE_DEVICE void + store_raw(const static_distributed_tensor& + dstr_tensor, + number = {}) const { - using vector_t = typename traits::vector_t; - using SFC_Ys = typename traits::SFC_Ys; + using vector_t = typename Base::Traits::vector_t; + using SFC_Ys = typename Base::Traits::SFC_Ys; - constexpr auto tile_dstr = TileDstr{}; + constexpr auto tile_dstr = typename Base::TileDstr{}; static constexpr bool oob_conditional_check = true; // loop over thread tensor space [y0, y1, ...] @@ -861,20 +690,21 @@ struct tile_window_linear // read from distributed tensor vector_t vec_value; - static_for<0, traits::ScalarPerVector, traits::PackedSize>{}([&](auto j) { + static_for<0, Base::Traits::ScalarPerVector, Base::Traits::PackedSize>{}([&](auto j) { constexpr auto idx_ys = generate_tuple( [&](auto jj) { - return jj == traits::VectorDimY ? (idx_ys_start[jj] + j) : idx_ys_start[jj]; + return jj == Base::Traits::VectorDimY ? (idx_ys_start[jj] + j) + : idx_ys_start[jj]; }, - number{}); + number{}); constexpr index_t d = tile_dstr.get_ys_to_d_descriptor().calculate_offset(idx_ys) / - traits::PackedSize; - vec_value.template get_as()(j / traits::PackedSize) = + Base::Traits::PackedSize; + vec_value.template get_as()(j / Base::Traits::PackedSize) = dstr_tensor.get_thread_buffer().template at(); }); // write into bottom tensor - get_bottom_tensor_view() + this->get_bottom_tensor_view() .template set_vectorized_elements_raw( bottom_tensor_thread_coord, linear_offset, bottom_tensor_flag, vec_value); }; @@ -883,15 +713,17 @@ struct tile_window_linear } template - CK_TILE_DEVICE void update(const static_distributed_tensor& dstr_tensor, - number = {}, - bool_constant = {}) const + CK_TILE_DEVICE void + update(const static_distributed_tensor& + dstr_tensor, + number = {}, + bool_constant = {}) const { - using vector_t = typename traits::vector_t; - using SFC_Ys = typename traits::SFC_Ys; + using vector_t = typename Base::Traits::vector_t; + using SFC_Ys = typename Base::Traits::SFC_Ys; - constexpr auto tile_dstr = TileDstr{}; + constexpr auto tile_dstr = typename Base::TileDstr{}; // loop over thread tensor space [y0, y1, ...] auto issue = [&](auto i_access_) { @@ -907,22 +739,24 @@ struct tile_window_linear // read from distributed tensor vector_t vec_value; - static_for<0, traits::ScalarPerVector, traits::PackedSize>{}([&](auto j) { + static_for<0, Base::Traits::ScalarPerVector, Base::Traits::PackedSize>{}([&](auto j) { constexpr auto idx_ys = generate_tuple( [&](auto jj) { - return jj == traits::VectorDimY ? (idx_ys_start[jj] + j) : idx_ys_start[jj]; + return jj == Base::Traits::VectorDimY ? (idx_ys_start[jj] + j) + : idx_ys_start[jj]; }, - number{}); + number{}); constexpr index_t d = tile_dstr.get_ys_to_d_descriptor().calculate_offset(idx_ys) / - traits::PackedSize; + Base::Traits::PackedSize; - vec_value.template get_as()(j / traits::PackedSize) = + vec_value.template get_as()( + j / Base::Traits::PackedSize) = dstr_tensor.get_thread_buffer().template at(); }); // write into bottom tensor - get_bottom_tensor_view().template update_vectorized_elements( + this->get_bottom_tensor_view().template update_vectorized_elements( bottom_tensor_thread_coord, linear_offset, bottom_tensor_flag, @@ -934,16 +768,18 @@ struct tile_window_linear } template - CK_TILE_DEVICE void update_raw(const static_distributed_tensor& dstr_tensor, - number = {}, - bool_constant = {}, - bool_constant = {}) const + CK_TILE_DEVICE void + update_raw(const static_distributed_tensor& + dstr_tensor, + number = {}, + bool_constant = {}, + bool_constant = {}) const { - using vector_t = typename traits::vector_t; - using SFC_Ys = typename traits::SFC_Ys; + using vector_t = typename Base::Traits::vector_t; + using SFC_Ys = typename Base::Traits::SFC_Ys; - constexpr auto tile_dstr = TileDstr{}; + constexpr auto tile_dstr = typename Base::TileDstr{}; // loop over thread tensor space [y0, y1, ...] auto issue = [&](auto i_access_) { @@ -959,22 +795,24 @@ struct tile_window_linear // read from distributed tensor vector_t vec_value; - static_for<0, traits::ScalarPerVector, traits::PackedSize>{}([&](auto j) { + static_for<0, Base::Traits::ScalarPerVector, Base::Traits::PackedSize>{}([&](auto j) { constexpr auto idx_ys = generate_tuple( [&](auto jj) { - return jj == traits::VectorDimY ? (idx_ys_start[jj] + j) : idx_ys_start[jj]; + return jj == Base::Traits::VectorDimY ? (idx_ys_start[jj] + j) + : idx_ys_start[jj]; }, - number{}); + number{}); constexpr index_t d = tile_dstr.get_ys_to_d_descriptor().calculate_offset(idx_ys) / - traits::PackedSize; + Base::Traits::PackedSize; - vec_value.template get_as()(j / traits::PackedSize) = + vec_value.template get_as()( + j / Base::Traits::PackedSize) = dstr_tensor.get_thread_buffer().template at(); }); // write into bottom tensor - get_bottom_tensor_view().template update_vectorized_elements_raw( + this->get_bottom_tensor_view().template update_vectorized_elements_raw( bottom_tensor_thread_coord, linear_offset, bottom_tensor_flag, @@ -985,14 +823,10 @@ struct tile_window_linear WINDOW_DISPATCH_ISSUE(); } - - // move thread's botom tensor coordiante - // [x0', x1', ... ] ==> [offset] - // also move window-origin - CK_TILE_DEVICE void move(const BottomTensorIndex& step) + // *_extended() functions acts like a virtual function with a default implementation exisiting + // in the base class + CK_TILE_DEVICE void move_extended(const typename Base::BottomTensorIndex& step) { - window_origin_ += step; - static_for<0, NumAccess, 1>{}([&](auto i_access) { constexpr auto IAccess = number{}; constexpr auto non_linear_id = number{}; @@ -1001,7 +835,7 @@ struct tile_window_linear if constexpr(need_update_non_linear_coord) { - move_tensor_coordinate(bottom_tensor_view_.get_tensor_descriptor(), + move_tensor_coordinate(this->bottom_tensor_view_.get_tensor_descriptor(), cached_coords_(non_linear_id), step); } @@ -1010,30 +844,29 @@ struct tile_window_linear auto tmp_coords = cached_coords_[non_linear_id]; constexpr auto linear_coord = get_bottom_linear_coordinate(IAccess); move_tensor_coordinate( - bottom_tensor_view_.get_tensor_descriptor(), tmp_coords, linear_coord); + this->bottom_tensor_view_.get_tensor_descriptor(), tmp_coords, linear_coord); cached_flags_(IAccess) = coordinate_has_valid_offset_assuming_top_index_is_valid( - bottom_tensor_view_.get_tensor_descriptor(), tmp_coords); + this->bottom_tensor_view_.get_tensor_descriptor(), tmp_coords); }); } - CK_TILE_DEVICE void set_window_origin(const BottomTensorIndex& new_window_origin) + CK_TILE_DEVICE void set_window_origin_extended(const typename Base::BottomTensorIndex&) { - window_origin_ = new_window_origin; - auto window_adaptor_thread_coord_tmp = make_tensor_adaptor_coordinate( - TileDstr{}.get_ps_ys_to_xs_adaptor(), - container_concat(make_tuple(get_warp_id(), get_lane_id()), - generate_tuple([&](auto) { return number<0>{}; }, number{}))); + typename Base::TileDstr{}.get_ps_ys_to_xs_adaptor(), + container_concat( + make_tuple(get_warp_id(), get_lane_id()), + generate_tuple([&](auto) { return number<0>{}; }, number{}))); - BottomTensorIndex bottom_tensor_thread_origin_idx_tmp = - window_origin_ + window_adaptor_thread_coord_tmp.get_bottom_index(); + typename Base::BottomTensorIndex bottom_tensor_thread_origin_idx_tmp = + this->window_origin_ + window_adaptor_thread_coord_tmp.get_bottom_index(); auto bottom_tensor_thread_coord_tmp = make_tensor_coordinate( - bottom_tensor_view_.get_tensor_descriptor(), bottom_tensor_thread_origin_idx_tmp); + this->bottom_tensor_view_.get_tensor_descriptor(), bottom_tensor_thread_origin_idx_tmp); // future load/store() calls (might allocate more registers) - using SFC_Ys = typename traits::SFC_Ys; + using SFC_Ys = typename Base::Traits::SFC_Ys; static_for<0, NumAccess, 1>{}([&](auto i_access) { constexpr auto non_linear_id = number{}; @@ -1049,10 +882,10 @@ struct tile_window_linear { constexpr auto idx_diff_ys = SFC_Ys::get_forward_step(i_access); // tuple of number constexpr auto idx_diff_ps_ys = container_concat( - generate_tuple([&](auto) { return number<0>{}; }, number{}), + generate_tuple([&](auto) { return number<0>{}; }, number{}), idx_diff_ys); - move_window_adaptor_and_bottom_tensor_thread_coordinate( + Base::move_window_adaptor_and_bottom_tensor_thread_coordinate( window_adaptor_thread_coord_tmp, bottom_tensor_thread_coord_tmp, idx_diff_ps_ys); @@ -1060,26 +893,9 @@ struct tile_window_linear }); } - CK_TILE_HOST_DEVICE void init_raw() { bottom_tensor_view_.init_raw(); } - - // this is the bottom tensor view - // [x0', x1', ...] ==> [offset] - BottomTensorView bottom_tensor_view_; - - // - WindowLengths window_lengths_; - - // origin ([x0', x1', ...]) of window on bottom tensor - BottomTensorIndex window_origin_; - - // Tile tensor distribution, which contains: - // 1. adaptor for window: [p0, p1, ..., y0, y1, ...] ==> [x0, x1, ...] - // 2. thread descriptor for thread tensor in register: [y0, y1, ...] ==> [d] - TileDstr tile_dstr_; - // this contains: - array cached_coords_; - array cached_flags_; + array cached_coords_; + array cached_flags_; }; #undef WINDOW_DISPATCH_ISSUE From 417a6b65b6436b533c64966a6b8825f12c31d8ef Mon Sep 17 00:00:00 2001 From: Adam Dickin Date: Thu, 22 May 2025 12:14:33 -0600 Subject: [PATCH 2/5] Add MIOPEN_REQ_LIBS_ONLY option for cmake to build only the libs MIOpen requires (#2224) * cut out anything we dont need for MIOpen to test * refactor exclusion code to be more streamlined. --- CMakeLists.txt | 15 +++++--- .../gpu/CMakeLists.txt | 38 +++++++++++++------ 2 files changed, 36 insertions(+), 17 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 13606975c0..2f1fc8892b 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -605,6 +605,9 @@ ENDIF() ENDFOREACH() add_custom_target(instances DEPENDS utility;${CK_DEVICE_INSTANCES} SOURCES ${INSTANCE_FILES}) + +option(MIOPEN_REQ_LIBS_ONLY "Build only the MIOpen required libraries" OFF) + add_subdirectory(library) if(NOT GPU_ARCHS AND USER_GPU_TARGETS) @@ -624,11 +627,13 @@ if(NOT GPU_ARCHS AND USER_GPU_TARGETS) endif() endif() -rocm_package_setup_component(profiler - LIBRARY_NAME composablekernel - PACKAGE_NAME ckprofiler -) -add_subdirectory(profiler) +if (NOT MIOPEN_REQ_LIBS_ONLY) + rocm_package_setup_component(profiler + LIBRARY_NAME composablekernel + PACKAGE_NAME ckprofiler + ) + add_subdirectory(profiler) +endif() if(CK_USE_CODEGEN AND (SUPPORTED_GPU_TARGETS MATCHES "gfx9" OR GPU_ARCHS)) add_subdirectory(codegen) diff --git a/library/src/tensor_operation_instance/gpu/CMakeLists.txt b/library/src/tensor_operation_instance/gpu/CMakeLists.txt index 97946207a1..73f3ae6cd6 100755 --- a/library/src/tensor_operation_instance/gpu/CMakeLists.txt +++ b/library/src/tensor_operation_instance/gpu/CMakeLists.txt @@ -103,6 +103,18 @@ function(add_instance_library INSTANCE_NAME) list(REMOVE_ITEM ARGN "${source}") endif() endforeach() + + if(MIOPEN_REQ_LIBS_ONLY) + message("Removing all sources that are not required for MIOpen") + foreach(source IN LISTS ARGN) + if(source MATCHES "gemm" OR + source MATCHES "mha" OR + source MATCHES "contraction" OR + source MATCHES "reduce") + list(REMOVE_ITEM ARGN "${source}") + endif() + endforeach() + endif() #message("remaining instances: ${ARGN}") #only continue if there are some source files left on the list if(ARGN) @@ -343,7 +355,7 @@ if(CK_DEVICE_OTHER_INSTANCES) DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/composable_kernel ) endif() -if(CK_DEVICE_GEMM_INSTANCES) +if(CK_DEVICE_GEMM_INSTANCES AND NOT MIOPEN_REQ_LIBS_ONLY) add_library(device_gemm_operations ${CK_DEVICE_GEMM_INSTANCES}) add_library(composablekernels::device_gemm_operations ALIAS device_gemm_operations) target_compile_features(device_gemm_operations PUBLIC) @@ -389,7 +401,7 @@ if(CK_DEVICE_CONV_INSTANCES) DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/composable_kernel ) endif() -if(CK_DEVICE_MHA_INSTANCES) +if(CK_DEVICE_MHA_INSTANCES AND NOT MIOPEN_REQ_LIBS_ONLY) set(gpu_list ${INST_TARGETS}) if(gpu_list MATCHES "gfx94" OR gpu_list MATCHES "gfx90a" OR gpu_list MATCHES "gfx95") add_library(device_mha_operations ${CK_DEVICE_MHA_INSTANCES}) @@ -411,7 +423,7 @@ if(CK_DEVICE_MHA_INSTANCES) ) endif() endif() -if(CK_DEVICE_CONTRACTION_INSTANCES) +if(CK_DEVICE_CONTRACTION_INSTANCES AND NOT MIOPEN_REQ_LIBS_ONLY) add_library(device_contraction_operations ${CK_DEVICE_CONTRACTION_INSTANCES}) add_library(composablekernels::device_contraction_operations ALIAS device_contraction_operations) target_compile_features(device_contraction_operations PUBLIC) @@ -433,7 +445,7 @@ if(CK_DEVICE_CONTRACTION_INSTANCES) DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/composable_kernel ) endif() -if(CK_DEVICE_REDUCTION_INSTANCES) +if(CK_DEVICE_REDUCTION_INSTANCES AND NOT MIOPEN_REQ_LIBS_ONLY) add_library(device_reduction_operations ${CK_DEVICE_REDUCTION_INSTANCES}) add_library(composablekernels::device_reduction_operations ALIAS device_reduction_operations) target_compile_features(device_reduction_operations PUBLIC) @@ -455,14 +467,16 @@ if(CK_DEVICE_REDUCTION_INSTANCES) ) endif() -add_library(device_operations INTERFACE) -target_link_libraries(device_operations INTERFACE - device_contraction_operations - device_conv_operations - device_gemm_operations - device_other_operations - device_reduction_operations - utility) +if(NOT MIOPEN_REQ_LIBS_ONLY) + add_library(device_operations INTERFACE) + target_link_libraries(device_operations INTERFACE + device_contraction_operations + device_conv_operations + device_gemm_operations + device_other_operations + device_reduction_operations + utility) +endif() set(DEV_OPS_INC_DIRS ${PROJECT_SOURCE_DIR}/include/ck/ From bc2551ac3b27edc31f20863e3a873508fb73aad2 Mon Sep 17 00:00:00 2001 From: Illia Silin <98187287+illsilin@users.noreply.github.com> Date: Thu, 22 May 2025 14:03:04 -0700 Subject: [PATCH 3/5] disable building device_mha_operations by default (#2225) --- CMakeLists.txt | 1 + library/src/tensor_operation_instance/gpu/CMakeLists.txt | 4 ++-- 2 files changed, 3 insertions(+), 2 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 2f1fc8892b..98698e5940 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -607,6 +607,7 @@ ENDFOREACH() add_custom_target(instances DEPENDS utility;${CK_DEVICE_INSTANCES} SOURCES ${INSTANCE_FILES}) option(MIOPEN_REQ_LIBS_ONLY "Build only the MIOpen required libraries" OFF) +option(BUILD_MHA_LIB "Build the static library for flash attention" OFF) add_subdirectory(library) diff --git a/library/src/tensor_operation_instance/gpu/CMakeLists.txt b/library/src/tensor_operation_instance/gpu/CMakeLists.txt index 73f3ae6cd6..aef40b8cb3 100755 --- a/library/src/tensor_operation_instance/gpu/CMakeLists.txt +++ b/library/src/tensor_operation_instance/gpu/CMakeLists.txt @@ -76,7 +76,7 @@ function(add_instance_library INSTANCE_NAME) endforeach() # Do not build mha instances if gfx94 or gfx90a targets are not on the target list foreach(source IN LISTS ARGN) - if(NOT INST_TARGETS MATCHES "gfx94" AND NOT INST_TARGETS MATCHES "gfx90a" AND NOT INST_TARGETS MATCHES "gfx95" AND source MATCHES "mha") + if((NOT BUILD_MHA_LIB OR (NOT INST_TARGETS MATCHES "gfx94" AND NOT INST_TARGETS MATCHES "gfx90a" AND NOT INST_TARGETS MATCHES "gfx95")) AND source MATCHES "mha") message("removing mha instance ${source} ") list(REMOVE_ITEM ARGN "${source}") endif() @@ -401,7 +401,7 @@ if(CK_DEVICE_CONV_INSTANCES) DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/composable_kernel ) endif() -if(CK_DEVICE_MHA_INSTANCES AND NOT MIOPEN_REQ_LIBS_ONLY) +if(CK_DEVICE_MHA_INSTANCES AND NOT MIOPEN_REQ_LIBS_ONLY AND BUILD_MHA_LIB) set(gpu_list ${INST_TARGETS}) if(gpu_list MATCHES "gfx94" OR gpu_list MATCHES "gfx90a" OR gpu_list MATCHES "gfx95") add_library(device_mha_operations ${CK_DEVICE_MHA_INSTANCES}) From 1b846143c669b7bcbba4fda5f9165bb270f88ea1 Mon Sep 17 00:00:00 2001 From: Illia Silin <98187287+illsilin@users.noreply.github.com> Date: Thu, 22 May 2025 15:41:17 -0700 Subject: [PATCH 4/5] Revert "Update the buffer load/store intrinsic names for clang>=20. (#2192)" (#2227) This reverts commit 58f9e9ffbc190188f85895deb952cb47cc89c403. --- .../ck_tile/18_flatmm/run_flatmm_example.inc | 74 +- .../amd_buffer_addressing_builtins.hpp | 20 +- include/ck_tile/core.hpp | 1 + .../arch/amd_buffer_addressing_builtins.hpp | 2559 +++++++++++++++++ include/ck_tile/core/tensor/buffer_view.hpp | 4 + 5 files changed, 2607 insertions(+), 51 deletions(-) create mode 100644 include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp diff --git a/example/ck_tile/18_flatmm/run_flatmm_example.inc b/example/ck_tile/18_flatmm/run_flatmm_example.inc index c191fff7d0..15a9df2c0c 100644 --- a/example/ck_tile/18_flatmm/run_flatmm_example.inc +++ b/example/ck_tile/18_flatmm/run_flatmm_example.inc @@ -4,22 +4,14 @@ #include template -constexpr const char* DataTypeToString() -{ - if constexpr(std::is_same_v) - { +constexpr const char* DataTypeToString() { + if constexpr (std::is_same_v) { return "fp16"; - } - else if constexpr(std::is_same_v) - { + } else if constexpr (std::is_same_v) { return "fp8"; - } - else if constexpr(std::is_same_v) - { + } else if constexpr (std::is_same_v) { return "bf8"; - } - else - { + } else { return "unknown"; } } @@ -120,9 +112,8 @@ float invoke_flatmm(ck_tile::DeviceMem& a_dev_buf, args.stride_B = stride_B; args.stride_C = stride_C; - float ave_time = - flatmm_calc( - args, ck_tile::stream_config{nullptr, true, 1, n_warmup, n_repeat}); + float ave_time = flatmm_calc( + args, ck_tile::stream_config{nullptr, true, 1, n_warmup, n_repeat}); std::size_t flop = std::size_t(2) * M * N * K; std::size_t num_byte = @@ -130,15 +121,18 @@ float invoke_flatmm(ck_tile::DeviceMem& a_dev_buf, float tflops = static_cast(flop) / 1.E9 / ave_time; float gb_per_sec = num_byte / 1.E6 / ave_time; - std::cout << "Run Flatmm kernel with DataType = " << DataTypeToString() - << " M =" << M << " N =" << N << " K =" << K << " StrideA =" << stride_A - << " StrideB =" << stride_B << " StrideC =" << stride_C << " : " << ave_time - << " ms, " << tflops << " TFlops, " << gb_per_sec << " GB/s, " << std::endl; + std::cout << "Run Flatmm kernel with DataType = " << DataTypeToString() << " M =" << M << " N =" << N << " K =" << K + << " StrideA =" << stride_A << " StrideB =" << stride_B << " StrideC =" << stride_C + << " : " << ave_time << " ms, " << tflops << " TFlops, " << gb_per_sec << " GB/s, " + << std::endl; return ave_time; } -template +template int run_flatmm_example_with_layouts(int argc, char* argv[], const ALayout a_layout = ALayout{}, @@ -153,7 +147,7 @@ int run_flatmm_example_with_layouts(int argc, using BDataType = typename GemmBasicTypeConfig::BDataType; using CDataType = typename GemmBasicTypeConfig::CDataType; using AccDataType = typename GemmBasicTypeConfig::AccDataType; - + ck_tile::index_t M = arg_parser.get_int("m"); ck_tile::index_t N = arg_parser.get_int("n"); ck_tile::index_t K = arg_parser.get_int("k"); @@ -188,7 +182,7 @@ int run_flatmm_example_with_layouts(int argc, c_rslt_host.SetZero(); // do pre-shuffle - std::string mfma = arg_parser.get_str("prec"); + std::string mfma = arg_parser.get_str("prec"); #if defined(USING_MFMA_16x16x32) && defined(ENABLE_FP8) ck_tile::index_t mfma_type = 1; #else @@ -199,18 +193,18 @@ int run_flatmm_example_with_layouts(int argc, b_shuffle_dev_buf.ToDevice(b_shuffle_host.data()); invoke_flatmm( - a_dev_buf, - b_shuffle_dev_buf, - c_dev_buf, - M, - N, - K, - stride_A, - stride_B, - stride_C, - kbatch, - n_warmup, - n_repeat); + a_dev_buf, + b_shuffle_dev_buf, + c_dev_buf, + M, + N, + K, + stride_A, + stride_B, + stride_C, + kbatch, + n_warmup, + n_repeat); c_dev_buf.FromDevice(c_rslt_host.data()); bool pass = true; @@ -225,9 +219,8 @@ int run_flatmm_example_with_layouts(int argc, a_host, b_origin_host, c_ref_host); const float max_accumulated_value = *std::max_element(c_ref_host.mData.begin(), c_ref_host.mData.end()); - const auto rtol_atol = calculate_rtol_atol( - K, kbatch, max_accumulated_value); - pass = ck_tile::check_err(c_rslt_host, + const auto rtol_atol = calculate_rtol_atol(K, kbatch, max_accumulated_value); + pass = ck_tile::check_err(c_rslt_host, c_ref_host, "Error: Incorrect results!", rtol_atol.at(ck_tile::number<0>{}), @@ -284,9 +277,8 @@ int run_flatmm_example_with_layouts(int argc, c_gpu_ref_dev_buf.FromDevice(c_gpu_ref_host.data()); const float max_accumulated_value = *std::max_element(c_gpu_ref_host.mData.begin(), c_gpu_ref_host.mData.end()); - const auto rtol_atol = calculate_rtol_atol( - K, kbatch, max_accumulated_value); - pass = ck_tile::check_err(c_rslt_host, + const auto rtol_atol = calculate_rtol_atol(K, kbatch, max_accumulated_value); + pass = ck_tile::check_err(c_rslt_host, c_gpu_ref_host, "Error: Incorrect results!", rtol_atol.at(ck_tile::number<0>{}), diff --git a/include/ck/utility/amd_buffer_addressing_builtins.hpp b/include/ck/utility/amd_buffer_addressing_builtins.hpp index 296c1d44d7..19869906dc 100644 --- a/include/ck/utility/amd_buffer_addressing_builtins.hpp +++ b/include/ck/utility/amd_buffer_addressing_builtins.hpp @@ -80,7 +80,7 @@ __device__ half2_t llvm_amdgcn_raw_buffer_atomic_add_fp16x2( int32x4_t rsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.fadd.v2f16"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.fadd.v2f16.v4i32"); // buffer atomic-add i32 __device__ int32_t llvm_amdgcn_raw_buffer_atomic_add_i32( @@ -88,7 +88,7 @@ __device__ int32_t llvm_amdgcn_raw_buffer_atomic_add_i32( int32x4_t rsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.add.i32"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.add.i32.v4i32"); // buffer atomic-add fp32 __device__ float llvm_amdgcn_raw_buffer_atomic_add_fp32( @@ -96,15 +96,15 @@ __device__ float llvm_amdgcn_raw_buffer_atomic_add_fp32( int32x4_t rsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.fadd.f32"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.fadd.f32.v4i32"); // buffer atomic-add fp32 -__device__ double -llvm_amdgcn_raw_buffer_atomic_max_fp64(double vdata, - int32x4_t rsrc, // dst_wave_buffer_resource - int voffset, // dst_thread_addr_offset - int soffset, // dst_wave_addr_offset - int glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.fmax.f64"); +__device__ double llvm_amdgcn_raw_buffer_atomic_max_fp64( + double vdata, + int32x4_t rsrc, // dst_wave_buffer_resource + int voffset, // dst_thread_addr_offset + int soffset, // dst_wave_addr_offset + int glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.fmax.f64.v4i32"); // memory coherency bit for buffer store/load instruction // check ISA manual for each GFX target @@ -827,7 +827,7 @@ llvm_amdgcn_raw_buffer_load_lds(int32x4_t rsrc, index_t voffset, index_t soffset, index_t offset, - index_t aux) __asm("llvm.amdgcn.raw.buffer.load.lds"); + index_t aux) __asm("llvm.amdgcn.raw.buffer.load.lds.v4i32"); #ifndef __HIPCC_RTC__ template diff --git a/include/ck_tile/core.hpp b/include/ck_tile/core.hpp index aa9411b2e1..27af59c192 100644 --- a/include/ck_tile/core.hpp +++ b/include/ck_tile/core.hpp @@ -9,6 +9,7 @@ #include "ck_tile/core/algorithm/space_filling_curve.hpp" #include "ck_tile/core/algorithm/static_encoding_pattern.hpp" #include "ck_tile/core/arch/amd_buffer_addressing.hpp" +#include "ck_tile/core/arch/amd_buffer_addressing_builtins.hpp" #include "ck_tile/core/arch/arch.hpp" #include "ck_tile/core/arch/generic_memory_space_atomic.hpp" #include "ck_tile/core/arch/utility.hpp" diff --git a/include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp b/include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp new file mode 100644 index 0000000000..0b9956cd01 --- /dev/null +++ b/include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp @@ -0,0 +1,2559 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved. + +#pragma once + +#if CK_TILE_USE_BUFFER_ADDRESSING_BUILTIN + +#include "ck_tile/core/numeric/integer.hpp" +#include "ck_tile/core/numeric/integral_constant.hpp" +#include "ck_tile/core/numeric/vector_type.hpp" +#include "ck_tile/core/container/container_helper.hpp" +#include "ck_tile/core/container/thread_buffer.hpp" +#include "ck_tile/core/utility/type_traits.hpp" +#include "ck_tile/core/utility/bit_cast.hpp" +#include "ck_tile/core/utility/functional.hpp" + +namespace ck_tile { + +// 128 bit SGPRs to supply buffer resource in buffer instructions +// https://rocm-documentation.readthedocs.io/en/latest/GCN_ISA_Manuals/testdocbook.html#vector-memory-buffer-instructions +struct __attribute__((packed)) buffer_resource +{ + const void* ptr; + uint32_t range; + uint32_t config; +}; + +CK_TILE_DEVICE int32x4_t make_wave_buffer_resource(const void* ptr, uint32_t size = 0xffffffff) +{ + buffer_resource res{ptr, size, CK_TILE_BUFFER_RESOURCE_3RD_DWORD}; + int32x4_t r = __builtin_bit_cast(int32x4_t, res); + r.x = __builtin_amdgcn_readfirstlane(r.x); + r.y = __builtin_amdgcn_readfirstlane(r.y); + r.z = __builtin_amdgcn_readfirstlane(r.z); + r.w = __builtin_amdgcn_readfirstlane(r.w); + return r; +} + +namespace impl { +// below type indicate the data type used for buffer load inline asm +// clang-format off +template struct buffer_load_trait; + +template struct buffer_load_trait<16, T> { using payload_t = fp32x4_t; }; +template struct buffer_load_trait<8 , T> { using payload_t = fp32x2_t; }; +template struct buffer_load_trait<4 , T> { using payload_t = float; }; +template struct buffer_load_trait<2 , T> { using payload_t = float; }; +template struct buffer_load_trait<1 , T> { using payload_t = float; }; + +#if CK_TILE_BUFFER_LOAD_RAW_BF16_WA +template<> struct buffer_load_trait<16, thread_buffer> { using payload_t = bf16x8_t; }; +template<> struct buffer_load_trait<8 , thread_buffer> { using payload_t = bf16x4_t; }; +template<> struct buffer_load_trait<4 , thread_buffer> { using payload_t = bf16x2_t; }; +#endif +// clang-format on +} // namespace impl + +// TODO: glc/slc/... +template +struct buffer_load; +#pragma clang diagnostic push +#pragma clang diagnostic ignored "-Wundefined-reinterpret-cast" +// TODO: strict aliasing rule seems fail when reinterpret_cast between vector type +// (exp_vector_type(xxx)) +template +struct buffer_load<16, pre_nop> +{ + template + CK_TILE_DEVICE void operator()(T& value, + int32x4_t res /*buffer resource*/, + index_t v_offset, + index_t /*s_offset*/, + index_t i_offset /*max 0xFFF*/, + index_t /*flag*/ = 0, + bool_constant = {}) + { + static_assert(sizeof(T) == 16); + using mbuf_t = typename impl::buffer_load_trait<16, T>::payload_t; + if constexpr(pre_nop) + asm volatile("s_nop 4\n" + "buffer_load_dwordx4 %0, %1, %2, 0 offen offset:%3" + : "+v"(reinterpret_cast(value)) + : "v"(v_offset), "s"(res), "n"(i_offset) + : "memory"); + else + asm volatile("buffer_load_dwordx4 %0, %1, %2, 0 offen offset:%3" + : "+v"(reinterpret_cast(value)) + : "v"(v_offset), "s"(res), "n"(i_offset) + : "memory"); + } +}; + +template +struct buffer_load<8, pre_nop> +{ + template + CK_TILE_DEVICE void operator()(T& value, + int32x4_t res /*buffer resource*/, + index_t v_offset, + index_t /*s_offset*/, + index_t i_offset /*max 0xFFF*/, + index_t /*flag*/ = 0, + bool_constant = {}) + { + static_assert(sizeof(T) == 8); + using mbuf_t = typename impl::buffer_load_trait<8, T>::payload_t; + if constexpr(pre_nop) + asm volatile("s_nop 4\n" + "buffer_load_dwordx2 %0, %1, %2, 0 offen offset:%3" + : "+v"(reinterpret_cast(value)) + : "v"(v_offset), "s"(res), "n"(i_offset) + : "memory"); + else + asm volatile("buffer_load_dwordx2 %0, %1, %2, 0 offen offset:%3" + : "+v"(reinterpret_cast(value)) + : "v"(v_offset), "s"(res), "n"(i_offset) + : "memory"); + } +}; + +template +struct buffer_load<4, pre_nop> +{ + template + CK_TILE_DEVICE void operator()(T& value, + int32x4_t res /*buffer resource*/, + index_t v_offset, + index_t /*s_offset*/, + index_t i_offset /*max 0xFFF*/, + index_t /*flag*/ = 0, + bool_constant = {}) + { + static_assert(sizeof(T) == 4); + using mbuf_t = typename impl::buffer_load_trait<4, T>::payload_t; + if constexpr(pre_nop) + asm volatile("s_nop 4\n" + "buffer_load_dword %0, %1, %2, 0 offen offset:%3" + : "+v"(reinterpret_cast(value)) + : "v"(v_offset), "s"(res), "n"(i_offset) + : "memory"); + else + asm volatile("buffer_load_dword %0, %1, %2, 0 offen offset:%3" + : "+v"(reinterpret_cast(value)) + : "v"(v_offset), "s"(res), "n"(i_offset) + : "memory"); + } +}; + +template +struct buffer_load<2, pre_nop> +{ + template + CK_TILE_DEVICE void operator()(T& value, + int32x4_t res /*buffer resource*/, + index_t v_offset, + index_t /*s_offset*/, + index_t i_offset /*max 0xFFF*/, + index_t /*flag*/ = 0, + bool_constant = {}) + { + static_assert(sizeof(T) == 4); // subdword is buggy, use dword buf and convert manually + using mbuf_t = typename impl::buffer_load_trait<2, T>::payload_t; + if constexpr(pre_nop) + asm volatile("s_nop 4\n" + "buffer_load_ushort %0, %1, %2, 0 offen offset:%3" + : "+v"(reinterpret_cast(value)) + : "v"(v_offset), "s"(res), "n"(i_offset) + : "memory"); + else + asm volatile("buffer_load_ushort %0, %1, %2, 0 offen offset:%3" + : "+v"(reinterpret_cast(value)) + : "v"(v_offset), "s"(res), "n"(i_offset) + : "memory"); + } +}; + +template +struct buffer_load<1, pre_nop> +{ + template + CK_TILE_DEVICE void operator()(T& value, + int32x4_t res /*buffer resource*/, + index_t v_offset, + index_t /*s_offset*/, + index_t i_offset /*max 0xFFF*/, + index_t /*flag*/ = 0, + bool_constant = {}) + { + static_assert(sizeof(T) == 4); + using mbuf_t = typename impl::buffer_load_trait<1, T>::payload_t; + if constexpr(pre_nop) + asm volatile("s_nop 4\n" + "buffer_load_ubyte %0, %1, %2, 0 offen offset:%3" + : "+v"(reinterpret_cast(value)) + : "v"(v_offset), "s"(res), "n"(i_offset) + : "memory"); + else + asm volatile("buffer_load_ubyte %0, %1, %2, 0 offen offset:%3" + : "+v"(reinterpret_cast(value)) + : "v"(v_offset), "s"(res), "n"(i_offset) + : "memory"); + } +}; + +template +struct buffer_load_if; + +template +struct buffer_load_if<16, pre_nop> +{ + template + CK_TILE_DEVICE void operator()(T& value, + int32x4_t res /*buffer resource*/, + index_t v_offset, + index_t /*s_offset*/, + index_t i_offset /*max 0xFFF*/, + index_t flag = 0, + bool_constant = {}) + { + static_assert(sizeof(T) == 16); + auto saved_exec = __builtin_amdgcn_read_exec(); + using mbuf_t = typename impl::buffer_load_trait<16, T>::payload_t; + static_assert(sizeof(mbuf_t) == sizeof(T)); + if constexpr(pre_nop) + asm volatile("s_nop 4\n" + "v_cmpx_le_u32 exec, 1, %4\n" + "buffer_load_dwordx4 %0, %1, %2, 0 offen offset:%3\n" + "s_mov_b64 exec %5" + : "+v"(reinterpret_cast(value)) + : "v"(v_offset), "s"(res), "n"(i_offset), "v"(flag), "s"(saved_exec) + : "memory"); + else + asm volatile("v_cmpx_le_u32 exec, 1, %4\n" + "buffer_load_dwordx4 %0, %1, %2, 0 offen offset:%3\n" + "s_mov_b64 exec %5" + : "+v"(reinterpret_cast(value)) + : "v"(v_offset), "s"(res), "n"(i_offset), "v"(flag), "s"(saved_exec) + : "memory"); + } +}; + +template +struct buffer_load_if<8, pre_nop> +{ + template + CK_TILE_DEVICE void operator()(T& value, + int32x4_t res /*buffer resource*/, + index_t v_offset, + index_t /*s_offset*/, + index_t i_offset /*max 0xFFF*/, + index_t flag = 0, + bool_constant = {}) + { + static_assert(sizeof(T) == 8); + auto saved_exec = __builtin_amdgcn_read_exec(); + using mbuf_t = typename impl::buffer_load_trait<8, T>::payload_t; + if constexpr(pre_nop) + asm volatile("s_nop 4\n" + "v_cmpx_le_u32 exec, 1, %4\n" + "buffer_load_dwordx2 %0, %1, %2, 0 offen offset:%3\n" + "s_mov_b64 exec %5" + : "+v"(reinterpret_cast(value)) + : "v"(v_offset), "s"(res), "n"(i_offset), "v"(flag), "s"(saved_exec) + : "memory"); + else + asm volatile("v_cmpx_le_u32 exec, 1, %4\n" + "buffer_load_dwordx2 %0, %1, %2, 0 offen offset:%3\n" + "s_mov_b64 exec %5" + : "+v"(reinterpret_cast(value)) + : "v"(v_offset), "s"(res), "n"(i_offset), "v"(flag), "s"(saved_exec) + : "memory"); + } +}; + +template +struct buffer_load_if<4, pre_nop> +{ + template + CK_TILE_DEVICE void operator()(T& value, + int32x4_t res /*buffer resource*/, + index_t v_offset, + index_t /*s_offset*/, + index_t i_offset /*max 0xFFF*/, + index_t flag = 0, + bool_constant = {}) + { + static_assert(sizeof(T) == 4); + auto saved_exec = __builtin_amdgcn_read_exec(); + using mbuf_t = typename impl::buffer_load_trait<4, T>::payload_t; + if constexpr(pre_nop) + asm volatile("s_nop 4\n" + "v_cmpx_le_u32 exec, 1, %4\n" + "buffer_load_dword %0, %1, %2, 0 offen offset:%3\n" + "s_mov_b64 exec %5" + : "+v"(reinterpret_cast(value)) + : "v"(v_offset), "s"(res), "n"(i_offset), "v"(flag), "s"(saved_exec) + : "memory"); + else + asm volatile("v_cmpx_le_u32 exec, 1, %4\n" + "buffer_load_dword %0, %1, %2, 0 offen offset:%3\n" + "s_mov_b64 exec %5" + : "+v"(reinterpret_cast(value)) + : "v"(v_offset), "s"(res), "n"(i_offset), "v"(flag), "s"(saved_exec) + : "memory"); + } +}; + +template +struct buffer_load_if<2, pre_nop> +{ + template + CK_TILE_DEVICE void operator()(T& value, + int32x4_t res /*buffer resource*/, + index_t v_offset, + index_t /*s_offset*/, + index_t i_offset /*max 0xFFF*/, + index_t flag = 0, + bool_constant = {}) + { + static_assert(sizeof(T) == 4); + auto saved_exec = __builtin_amdgcn_read_exec(); + using mbuf_t = typename impl::buffer_load_trait<2, T>::payload_t; + if constexpr(pre_nop) + asm volatile("s_nop 4\n" + "v_cmpx_le_u32 exec, 1, %4\n" + "buffer_load_ushort %0, %1, %2, 0 offen offset:%3\n" + "s_mov_b64 exec %5" + : "+v"(reinterpret_cast(value)) + : "v"(v_offset), "s"(res), "n"(i_offset), "v"(flag), "s"(saved_exec) + : "memory"); + else + asm volatile("v_cmpx_le_u32 exec, 1, %4\n" + "buffer_load_ushort %0, %1, %2, 0 offen offset:%3\n" + "s_mov_b64 exec %5" + : "+v"(reinterpret_cast(value)) + : "v"(v_offset), "s"(res), "n"(i_offset), "v"(flag), "s"(saved_exec) + : "memory"); + } +}; + +template +struct buffer_load_if<1, pre_nop> +{ + template + CK_TILE_DEVICE void operator()(T& value, + int32x4_t res /*buffer resource*/, + index_t v_offset, + index_t /*s_offset*/, + index_t i_offset /*max 0xFFF*/, + index_t flag = 0, + bool_constant = {}) + { + static_assert(sizeof(T) == 4); + auto saved_exec = __builtin_amdgcn_read_exec(); + using mbuf_t = typename impl::buffer_load_trait<1, T>::payload_t; + if constexpr(pre_nop) + asm volatile("s_nop 4\n" + "v_cmpx_le_u32 exec, 1, %4\n" + "buffer_load_ubyte %0, %1, %2, 0 offen offset:%3\n" + "s_mov_b64 exec %5" + : "+v"(reinterpret_cast(value)) + : "v"(v_offset), "s"(res), "n"(i_offset), "v"(flag), "s"(saved_exec) + : "memory"); + else + asm volatile("v_cmpx_le_u32 exec, 1, %4\n" + "buffer_load_ubyte %0, %1, %2, 0 offen offset:%3\n" + "s_mov_b64 exec %5" + : "+v"(reinterpret_cast(value)) + : "v"(v_offset), "s"(res), "n"(i_offset), "v"(flag), "s"(saved_exec) + : "memory"); + } +}; +#pragma clang diagnostic pop // "-Wundefined-reinterpret-cast" +template +struct buffer_store; + +template <> +struct buffer_store<16> +{ + template + CK_TILE_DEVICE void operator()(const T& value, + int32x4_t res /*buffer resource*/, + index_t v_offset, + index_t /*s_offset*/, + index_t i_offset /*max 0xFFF*/, + index_t /*flag*/ = 1) + { + static_assert(sizeof(T) == 16); + using mbuf_t = fp32x4_t; + asm volatile("buffer_store_dwordx4 %0, %1, %2, 0 offen offset:%3" + : + : "v"(bit_cast(value)), "v"(v_offset), "s"(res), "n"(i_offset) + : "memory"); + } +}; + +template <> +struct buffer_store<8> +{ + template + CK_TILE_DEVICE void operator()(const T& value, + int32x4_t res /*buffer resource*/, + index_t v_offset, + index_t /*s_offset*/, + index_t i_offset /*max 0xFFF*/, + index_t /*flag*/ = 1) + { + static_assert(sizeof(T) == 8); + using mbuf_t = fp32x2_t; + asm volatile("buffer_store_dwordx2 %0, %1, %2, 0 offen offset:%3" + : + : "v"(bit_cast(value)), "v"(v_offset), "s"(res), "n"(i_offset) + : "memory"); + } +}; + +template <> +struct buffer_store<4> +{ + template + CK_TILE_DEVICE void operator()(const T& value, + int32x4_t res /*buffer resource*/, + index_t v_offset, + index_t /*s_offset*/, + index_t i_offset /*max 0xFFF*/, + index_t /*flag*/ = 1) + { + static_assert(sizeof(T) == 4); + using mbuf_t = float; + asm volatile("buffer_store_dword %0, %1, %2, 0 offen offset:%3" + : + : "v"(bit_cast(value)), "v"(v_offset), "s"(res), "n"(i_offset) + : "memory"); + } +}; + +template <> +struct buffer_store<2> +{ + template + CK_TILE_DEVICE void operator()(const T& value, + int32x4_t res /*buffer resource*/, + index_t v_offset, + index_t /*s_offset*/, + index_t i_offset /*max 0xFFF*/, + index_t /*flag*/ = 1) + { + static_assert(sizeof(T) == 2); + using mbuf_t = short; + asm volatile("buffer_store_short %0, %1, %2, 0 offen offset:%3" + : + : "v"(bit_cast(value)), "v"(v_offset), "s"(res), "n"(i_offset) + : "memory"); + } +}; + +template <> +struct buffer_store<1> +{ + template + CK_TILE_DEVICE void operator()(const T& value, + int32x4_t res /*buffer resource*/, + index_t v_offset, + index_t /*s_offset*/, + index_t i_offset /*max 0xFFF*/, + index_t /*flag*/ = 1) + { + static_assert(sizeof(T) == 4); + using mbuf_t = float; + asm volatile("buffer_store_byte %0, %1, %2, 0 offen offset:%3" + : + : "v"(bit_cast(value)), "v"(v_offset), "s"(res), "n"(i_offset) + : "memory"); + } +}; + +template +struct buffer_store_if; + +template <> +struct buffer_store_if<16> +{ + template + CK_TILE_DEVICE void operator()(const T& value, + int32x4_t res /*buffer resource*/, + index_t v_offset, + index_t /*s_offset*/, + index_t i_offset /*max 0xFFF*/, + index_t flag = 1) + { + static_assert(sizeof(T) == 16); + auto save_exec = __builtin_amdgcn_read_exec(); + using mbuf_t = fp32x4_t; + asm volatile("v_cmpx_le_u32 exec, 1, %4\n" + "buffer_store_dwordx4 %0, %1, %2, 0 offen offset:%3\n" + "s_mov_b64 exec %5" + : + : "v"(bit_cast(value)), + "v"(v_offset), + "s"(res), + "n"(i_offset), + "v"(flag), + "s"(save_exec) + : "memory"); + } +}; + +template <> +struct buffer_store_if<8> +{ + template + CK_TILE_DEVICE void operator()(const T& value, + int32x4_t res /*buffer resource*/, + index_t v_offset, + index_t /*s_offset*/, + index_t i_offset /*max 0xFFF*/, + index_t flag = 1) + { + static_assert(sizeof(T) == 8); + auto save_exec = __builtin_amdgcn_read_exec(); + // TODO: ugly. rocm-6.0/6.1 seems neet bit_cast to same base type to avoid scratch + using mbuf_t = ext_vector_t; + asm volatile("v_cmpx_le_u32 exec, 1, %4\n" + "buffer_store_dwordx2 %0, %1, %2, 0 offen offset:%3\n" + "s_mov_b64 exec %5" + : + : "v"(bit_cast(value)), + "v"(v_offset), + "s"(res), + "n"(i_offset), + "v"(flag), + "s"(save_exec) + : "memory"); + } +}; + +template <> +struct buffer_store_if<4> +{ + template + CK_TILE_DEVICE void operator()(const T& value, + int32x4_t res /*buffer resource*/, + index_t v_offset, + index_t /*s_offset*/, + index_t i_offset /*max 0xFFF*/, + index_t flag = 1) + { + static_assert(sizeof(T) == 4); + auto save_exec = __builtin_amdgcn_read_exec(); + using mbuf_t = float; + asm volatile("v_cmpx_le_u32 exec, 1, %4\n" + "buffer_store_dword %0, %1, %2, 0 offen offset:%3\n" + "s_mov_b64 exec %5" + : + : "v"(bit_cast(value)), + "v"(v_offset), + "s"(res), + "n"(i_offset), + "v"(flag), + "s"(save_exec) + : "memory"); + } +}; + +template <> +struct buffer_store_if<2> +{ + template + CK_TILE_DEVICE void operator()(const T& value, + int32x4_t res /*buffer resource*/, + index_t v_offset, + index_t /*s_offset*/, + index_t i_offset /*max 0xFFF*/, + index_t flag = 1) + { + static_assert(sizeof(T) == 2); + auto save_exec = __builtin_amdgcn_read_exec(); + using mbuf_t = short; + asm volatile("v_cmpx_le_u32 exec, 1, %4\n" + "buffer_store_short %0, %1, %2, 0 offen offset:%3\n" + "s_mov_b64 exec %5" + : + : "v"(bit_cast(value)), + "v"(v_offset), + "s"(res), + "n"(i_offset), + "v"(flag), + "s"(save_exec) + : "memory"); + } +}; + +template <> +struct buffer_store_if<1> +{ + template + CK_TILE_DEVICE void operator()(const T& value, + int32x4_t res /*buffer resource*/, + index_t v_offset, + index_t /*s_offset*/, + index_t i_offset /*max 0xFFF*/, + index_t flag = 1) + { + static_assert(sizeof(T) == 4); + auto save_exec = __builtin_amdgcn_read_exec(); + using mbuf_t = float; + asm volatile("v_cmpx_le_u32 exec, 1, %4\n" + "buffer_store_byte %0, %1, %2, 0 offen offset:%3\n" + "s_mov_b64 exec %5" + : + : "v"(bit_cast(value)), + "v"(v_offset), + "s"(res), + "n"(i_offset), + "v"(flag), + "s"(save_exec) + : "memory"); + } +}; + +CK_TILE_DEVICE void buffer_load_fence(index_t cnt = 0) +{ + asm volatile("s_waitcnt vmcnt(%0)" : : "n"(cnt) : "memory"); +} + +CK_TILE_DEVICE void lds_load_fence(index_t cnt = 0) +{ + asm volatile("s_waitcnt lgkmcnt(%0)" : : "n"(cnt) : "memory"); +} + +template +struct buffer_atomic_add_if; + +template +struct buffer_atomic_add_if +{ + template + CK_TILE_DEVICE void operator()(const T& value, + int32x4_t res /*buffer resource*/, + index_t v_offset, + index_t /*s_offset*/, + index_t i_offset /*max 0xFFF*/, + index_t flag = 1) + { + static_assert(sizeof(T) == 4); + auto save_exec = __builtin_amdgcn_read_exec(); + using mbuf_t = float; + asm volatile("v_cmpx_le_u32 exec, 1, %4\n" + "global_atomic_pk_add_bf16 %0, %1, %2 offset:%3\n" + "s_mov_b64 exec %5" + : + : "v"(v_offset), + "v"(bit_cast(value)), + "s"(res.xy), + "n"(i_offset), + "v"(flag), + "s"(save_exec) + : "memory"); + } +}; + +template +struct buffer_atomic_add; + +template +struct buffer_atomic_add +{ + template + CK_TILE_DEVICE void operator()(const T& value, + int32x4_t res /*buffer resource*/, + index_t v_offset, + index_t /*s_offset*/, + index_t i_offset /*max 0xFFF*/, + index_t /*flag = 1*/) + { + static_assert(sizeof(T) == 4); + using mbuf_t = float; + asm volatile("global_atomic_pk_add_bf16 %0, %1, %2 offset:%3" + : + : "v"(v_offset), "v"(bit_cast(value)), "s"(res.xy), "n"(i_offset) + : "memory"); + } +}; + +namespace impl { +// below type indicate the data type used for buffer load inline asm +// clang-format off +template struct smem_load_trait; + +template struct smem_load_trait<16, T> { using payload_t = fp32x4_t; }; +template struct smem_load_trait<8 , T> { using payload_t = fp32x2_t; }; +template struct smem_load_trait<4 , T> { using payload_t = float; }; +template struct smem_load_trait<2 , T> { using payload_t = float; }; +template struct smem_load_trait<1 , T> { using payload_t = float; }; + +// clang-format on +} // namespace impl + +// NOTE: smem load/store no need pre_nop to make sure dependency by sw, happy :) +template +struct smem_load; + +template <> +struct smem_load<16> +{ + template + CK_TILE_DEVICE void operator()(T& value, index_t v_offset, index_t i_offset) + { + static_assert(sizeof(T) == 16); + using mbuf_t = typename impl::smem_load_trait<16, T>::payload_t; + asm volatile("ds_read_b128 %0, %1 offset:%2" + : "=v"(reinterpret_cast(value)) // ! direct write + : "v"(v_offset), "n"(i_offset) + : "memory"); + } +}; + +template <> +struct smem_load<8> +{ + template + CK_TILE_DEVICE void operator()(T& value, index_t v_offset, index_t i_offset) + { + static_assert(sizeof(T) == 8); + using mbuf_t = typename impl::smem_load_trait<8, T>::payload_t; + asm volatile("ds_read_b64 %0, %1 offset:%2" + : "=v"(reinterpret_cast(value)) // ! direct write + : "v"(v_offset), "n"(i_offset) + : "memory"); + } +}; + +template <> +struct smem_load<4> +{ + template + CK_TILE_DEVICE void operator()(T& value, index_t v_offset, index_t i_offset) + { + static_assert(sizeof(T) == 4); + using mbuf_t = typename impl::smem_load_trait<4, T>::payload_t; + asm volatile("ds_read_b32 %0, %1 offset:%2" + : "=v"(reinterpret_cast(value)) // ! direct write + : "v"(v_offset), "n"(i_offset) + : "memory"); + } +}; + +template <> +struct smem_load<2> +{ + template + CK_TILE_DEVICE void operator()(T& value, index_t v_offset, index_t i_offset) + { + static_assert(sizeof(T) == 4); // subdword is buggy, use dword buf and convert manually + using mbuf_t = typename impl::smem_load_trait<1, T>::payload_t; + asm volatile("ds_read_u16 %0, %1 offset:%2" + : "=v"(reinterpret_cast(value)) // ! direct write + : "v"(v_offset), "n"(i_offset) + : "memory"); + } +}; + +template <> +struct smem_load<1> +{ + template + CK_TILE_DEVICE void operator()(T& value, index_t v_offset, index_t i_offset) + { + static_assert(sizeof(T) == 4); + using mbuf_t = typename impl::smem_load_trait<1, T>::payload_t; + asm volatile("ds_read_u8 %0, %1 offset:%2" + : "=v"(reinterpret_cast(value)) // ! direct write + : "v"(v_offset), "n"(i_offset) + : "memory"); + } +}; + +// clang-format off +namespace impl{ + +// can't use "+v" since there could be potential extra move(read/write) +// use "v" can help remove such duplicated moves +// besides, fake this as "memory" operation to force later valu after this fence +// TODO: may have scratch (because this is memory?) +// need to reduce extra move inside compiler +template +CK_TILE_DEVICE void insert_dummy_dep_per_dword(array& b) +{ + constexpr auto kSize = remove_cvref_t::size(); + static_for<0, kSize, 1>{}([&](auto i){ + asm volatile(" " : : "v"(b.get(number{})) : "memory"); + }); +} +#if 1 +// below specialization just merge size() of dwords into single section +template<> +CK_TILE_DEVICE void insert_dummy_dep_per_dword<2>(array& b) +{ + asm volatile(" " : : "v"(b.get(number<0>{})), "v"(b.get(number<1>{})) : "memory"); +} + +template<> +CK_TILE_DEVICE void insert_dummy_dep_per_dword<3>(array& b) +{ + asm volatile(" " : : "v"(b.get(number<0>{})), "v"(b.get(number<1>{})), "v"(b.get(number<2>{})) : "memory"); +} + +template<> +CK_TILE_DEVICE void insert_dummy_dep_per_dword<4>(array& b) +{ + asm volatile(" " : : "v"(b.get(number<0>{})), "v"(b.get(number<1>{})), "v"(b.get(number<2>{})), "v"(b.get(number<3>{})) : "memory"); +} + +template<> +CK_TILE_DEVICE void insert_dummy_dep_per_dword<8>(array& b) +{ + asm volatile(" " : : "v"(b.get(number<0>{})), "v"(b.get(number<1>{})), "v"(b.get(number<2>{})), "v"(b.get(number<3>{})), + "v"(b.get(number<4>{})), "v"(b.get(number<5>{})), "v"(b.get(number<6>{})), "v"(b.get(number<7>{})) : "memory"); +} + +template<> +CK_TILE_DEVICE void insert_dummy_dep_per_dword<16>(array& b) +{ + asm volatile(" " : : "v"(b.get(number<0>{})), "v"(b.get(number<1>{})), "v"(b.get(number<2>{})), "v"(b.get(number<3>{})), + "v"(b.get(number<4>{})), "v"(b.get(number<5>{})), "v"(b.get(number<6>{})), "v"(b.get(number<7>{})), + "v"(b.get(number<8>{})), "v"(b.get(number<9>{})), "v"(b.get(number<10>{})), "v"(b.get(number<11>{})), + "v"(b.get(number<12>{})), "v"(b.get(number<13>{})), "v"(b.get(number<14>{})), "v"(b.get(number<15>{})) : "memory"); +} + +template<> +CK_TILE_DEVICE void insert_dummy_dep_per_dword<32>(array& b) +{ + asm volatile(" " : : "v"(b.get(number<0>{})), "v"(b.get(number<1>{})), "v"(b.get(number<2>{})), "v"(b.get(number<3>{})), + "v"(b.get(number<4>{})), "v"(b.get(number<5>{})), "v"(b.get(number<6>{})), "v"(b.get(number<7>{})), + "v"(b.get(number<8>{})), "v"(b.get(number<9>{})), "v"(b.get(number<10>{})), "v"(b.get(number<11>{})), + "v"(b.get(number<12>{})), "v"(b.get(number<13>{})), "v"(b.get(number<14>{})), "v"(b.get(number<15>{})), + "v"(b.get(number<16>{})), "v"(b.get(number<17>{})), "v"(b.get(number<18>{})), "v"(b.get(number<19>{})), + "v"(b.get(number<20>{})), "v"(b.get(number<21>{})), "v"(b.get(number<22>{})), "v"(b.get(number<23>{})), + "v"(b.get(number<24>{})), "v"(b.get(number<25>{})), "v"(b.get(number<26>{})), "v"(b.get(number<27>{})), + "v"(b.get(number<28>{})), "v"(b.get(number<29>{})), "v"(b.get(number<30>{})), "v"(b.get(number<31>{})) : "memory"); +} +#endif +CK_TILE_DEVICE void insert_dummy_dep() {} + +template +CK_TILE_DEVICE void insert_dummy_dep(T & buffer) +{ + // TODO: indeed we expect T to be multiple of dword. subdword is always buggy + using da_type = array; + auto & dummy = reinterpret_cast(buffer); + insert_dummy_dep_per_dword(dummy); +} + +template +CK_TILE_DEVICE void insert_dummy_dep(Tx& bx, Ty&... by) +{ + insert_dummy_dep(bx); + insert_dummy_dep(by...); +} +} +// clang-format on +template +CK_TILE_DEVICE void buffer_load_fence(index_t cnt = 0, T&... o) +{ + asm volatile("s_waitcnt vmcnt(%0)" : : "n"(cnt) : "memory"); + impl::insert_dummy_dep(o...); +} + +CK_TILE_DEVICE void buffer_store_fence(index_t cnt = 0) +{ + asm volatile("s_waitcnt vmcnt(%0)" : : "n"(cnt) : "memory"); +} + +CK_TILE_DEVICE auto async_load_fence_raw(index_t cnt = 0) +{ + asm volatile("s_waitcnt vmcnt(%0)" : : "n"(cnt) : "memory"); +} + +// buffer load i8 +CK_TILE_DEVICE_EXTERN int8_t +llvm_amdgcn_raw_buffer_load_i8(int32x4_t srsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.i8.v4i32"); + +CK_TILE_DEVICE_EXTERN int8x2_t +llvm_amdgcn_raw_buffer_load_i8x2(int32x4_t srsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v2i8.v4i32"); + +CK_TILE_DEVICE_EXTERN int8x4_t +llvm_amdgcn_raw_buffer_load_i8x4(int32x4_t srsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v4i8.v4i32"); + +// buffer load i16 +CK_TILE_DEVICE_EXTERN int16_t +llvm_amdgcn_raw_buffer_load_i16(int32x4_t srsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.i16.v4i32"); + +CK_TILE_DEVICE_EXTERN int16x2_t +llvm_amdgcn_raw_buffer_load_i16x2(int32x4_t srsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v2i16.v4i32"); + +CK_TILE_DEVICE_EXTERN int16x4_t +llvm_amdgcn_raw_buffer_load_i16x4(int32x4_t srsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v4i16.v4i32"); + +// buffer load i32 +CK_TILE_DEVICE_EXTERN int32_t +llvm_amdgcn_raw_buffer_load_i32(int32x4_t srsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.i32.v4i32"); + +CK_TILE_DEVICE_EXTERN int32x2_t +llvm_amdgcn_raw_buffer_load_i32x2(int32x4_t srsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v2i32.v4i32"); + +CK_TILE_DEVICE_EXTERN int32x4_t +llvm_amdgcn_raw_buffer_load_i32x4(int32x4_t srsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v4i32.v4i32"); + +// buffer load fp16 +CK_TILE_DEVICE_EXTERN _Float16 +llvm_amdgcn_raw_buffer_load_fp16(int32x4_t srsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.f16.v4i32"); + +CK_TILE_DEVICE_EXTERN fp16x2_t llvm_amdgcn_raw_buffer_load_fp16x2( + int32x4_t srsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v2f16.v4i32"); + +CK_TILE_DEVICE_EXTERN fp16x4_t llvm_amdgcn_raw_buffer_load_fp16x4( + int32x4_t srsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v4f16.v4i32"); + +// buffer load fp32 +CK_TILE_DEVICE_EXTERN float +llvm_amdgcn_raw_buffer_load_fp32(int32x4_t srsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.f32.v4i32"); + +CK_TILE_DEVICE_EXTERN fp32x2_t llvm_amdgcn_raw_buffer_load_fp32x2( + int32x4_t srsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v2f32.v4i32"); + +CK_TILE_DEVICE_EXTERN fp32x4_t llvm_amdgcn_raw_buffer_load_fp32x4( + int32x4_t srsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v4f32.v4i32"); + +// buffer store i8 +CK_TILE_DEVICE_EXTERN void +llvm_amdgcn_raw_buffer_store_i8(int8_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.i8.v4i32"); + +CK_TILE_DEVICE_EXTERN void +llvm_amdgcn_raw_buffer_store_i8x2(int8x2_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2i8.v4i32"); + +CK_TILE_DEVICE_EXTERN void +llvm_amdgcn_raw_buffer_store_i8x4(int8x4_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4i8.v4i32"); + +// buffer store i16 +CK_TILE_DEVICE_EXTERN void +llvm_amdgcn_raw_buffer_store_i16(int16_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.i16.v4i32"); + +CK_TILE_DEVICE_EXTERN void llvm_amdgcn_raw_buffer_store_i16x2( + int16x2_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2i16.v4i32"); + +CK_TILE_DEVICE_EXTERN void llvm_amdgcn_raw_buffer_store_i16x4( + int16x4_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4i16.v4i32"); + +// buffer store i32 +CK_TILE_DEVICE_EXTERN void +llvm_amdgcn_raw_buffer_store_i32(int32_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.i32.v4i32"); + +// buffer store ui16 +CK_TILE_DEVICE_EXTERN void +llvm_amdgcn_raw_buffer_store_ui16(uint16_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.i16.v4i32"); + +CK_TILE_DEVICE_EXTERN void llvm_amdgcn_raw_buffer_store_ui16x2( + uint16x2_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2i16.v4i32"); + +CK_TILE_DEVICE_EXTERN void llvm_amdgcn_raw_buffer_store_ui16x4( + uint16x4_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4i16.v4i32"); + +CK_TILE_DEVICE_EXTERN void llvm_amdgcn_raw_buffer_store_i32x2( + int32x2_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2i32.v4i32"); + +CK_TILE_DEVICE_EXTERN void llvm_amdgcn_raw_buffer_store_i32x4( + int32x4_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4i32.v4i32"); + +// buffer store fp16 +CK_TILE_DEVICE_EXTERN void +llvm_amdgcn_raw_buffer_store_fp16(_Float16 vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.f16.v4i32"); + +CK_TILE_DEVICE_EXTERN void llvm_amdgcn_raw_buffer_store_fp16x2( + fp16x2_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2f16.v4i32"); + +CK_TILE_DEVICE_EXTERN void llvm_amdgcn_raw_buffer_store_fp16x4( + fp16x4_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4f16.v4i32"); + +// buffer store fp32 +CK_TILE_DEVICE_EXTERN void +llvm_amdgcn_raw_buffer_store_fp32(float vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.f32.v4i32"); + +CK_TILE_DEVICE_EXTERN void llvm_amdgcn_raw_buffer_store_fp32x2( + fp32x2_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2f32.v4i32"); + +CK_TILE_DEVICE_EXTERN void llvm_amdgcn_raw_buffer_store_fp32x4( + fp32x4_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4f32.v4i32"); + +// buffer atomic-add fp16 +CK_TILE_DEVICE_EXTERN fp16x2_t llvm_amdgcn_raw_buffer_atomic_add_fp16x2( + fp16x2_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.fadd.v2f16.v4i32"); + +// buffer atomic-add i32 +CK_TILE_DEVICE_EXTERN int32_t llvm_amdgcn_raw_buffer_atomic_add_i32( + int32_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.add.i32.v4i32"); + +// buffer atomic-add fp32 +CK_TILE_DEVICE_EXTERN float llvm_amdgcn_raw_buffer_atomic_add_fp32( + float vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.fadd.f32.v4i32"); + +// buffer atomic-max fp64 +CK_TILE_DEVICE_EXTERN double llvm_amdgcn_raw_buffer_atomic_max_fp64( + double vdata, + int32x4_t rsrc, // dst_wave_buffer_resource + int voffset, // dst_thread_addr_offset + int soffset, // dst_wave_addr_offset + int glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.fmax.f64.v4i32"); + +// Direct loads from global to LDS. +CK_TILE_DEVICE_EXTERN void +llvm_amdgcn_raw_buffer_load_lds(int32x4_t rsrc, + __attribute__((address_space(3))) uint32_t* lds_ptr, + index_t size, + index_t voffset, + index_t soffset, + index_t offset, + index_t aux) __asm("llvm.amdgcn.raw.buffer.load.lds.v4i32"); + +template +CK_TILE_DEVICE void async_buffer_load_dword_v(void* smem, + int32x4_t rsrc, + index_t voffset, + index_t /*soffset*/, + index_t ioffset /*max 0xFFF*/, + index_t /*flag*/ = 0, + bool_constant = {}) +{ + if constexpr(pre_nop) + asm volatile("s_nop 4\n" + "buffer_load_dword %1, %2, 0 offen offset:%3 lds" + : "=r"(smem) /*dummy dependency for smem*/ + : "v"(voffset), "s"(rsrc), "n"(ioffset) + : "memory"); + else + asm volatile("buffer_load_dword %1, %2, 0 offen offset:%3 lds" + : "=r"(smem) /*dummy dependency for smem*/ + : "v"(voffset), "s"(rsrc), "n"(ioffset) + : "memory"); +} + +CK_TILE_DEVICE void async_buffer_load_fence(index_t cnt = 0) +{ + asm volatile("s_waitcnt vmcnt(%0)" : : "n"(cnt) : "memory"); +} + +// memory coherency bit for buffer store/load instruction +// check ISA manual for each GFX target +// e.g. for +// https://www.amd.com/system/files/TechDocs/instinct-mi200-cdna2-instruction-set-architecture.pdf, +// page 67~68 +enum struct amd_buffer_coherence_enum +{ + coherence_default = 0, // default value + glc = 1, + slc = 2, + glc_slc = 3, +}; + +template +CK_TILE_DEVICE thread_buffer +amd_buffer_load_impl_with_bytes(int32x4_t src_wave_buffer_resource, + index_t src_thread_addr_offset, + index_t src_wave_addr_offset) +{ + static_assert(N == 1 || N == 2 || N == 4 || N == 8 || N == 16 || N == 32 || N == 64, + "wrong! not implemented"); + + using rtn_type = thread_buffer; + + if constexpr(N == 1) + { + return bit_cast(llvm_amdgcn_raw_buffer_load_i8(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset, + static_cast(coherence))); + } + else if constexpr(N == 2) + { + + int16_t tmp = llvm_amdgcn_raw_buffer_load_i16(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset, + static_cast(coherence)); + + return bit_cast(tmp); + } + else if constexpr(N == 4) + { + int32_t tmp = llvm_amdgcn_raw_buffer_load_i32(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset, + static_cast(coherence)); + + return bit_cast(tmp); + } + else if constexpr(N == 8) + { + int32x2_t tmp = llvm_amdgcn_raw_buffer_load_i32x2(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset, + static_cast(coherence)); + + return bit_cast(tmp); + } + else if constexpr(N == 16) + { + int32x4_t tmp = llvm_amdgcn_raw_buffer_load_i32x4(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset, + static_cast(coherence)); + return bit_cast(tmp); + } + else if constexpr(N == 32) + { + int32x4_t tmp0 = llvm_amdgcn_raw_buffer_load_i32x4(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset, + static_cast(coherence)); + int32x4_t tmp1 = + llvm_amdgcn_raw_buffer_load_i32x4(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset + 4 * sizeof(int32_t), + static_cast(coherence)); + thread_buffer tmp; + + tmp.template get_as()(number<0>{}) = tmp0; + tmp.template get_as()(number<1>{}) = tmp1; + + return bit_cast(tmp); + } + else if constexpr(N == 64) + { + int32x4_t tmp0 = llvm_amdgcn_raw_buffer_load_i32x4(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset, + static_cast(coherence)); + int32x4_t tmp1 = + llvm_amdgcn_raw_buffer_load_i32x4(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset + 4 * sizeof(int32_t), + static_cast(coherence)); + int32x4_t tmp2 = + llvm_amdgcn_raw_buffer_load_i32x4(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset + 8 * sizeof(int32_t), + static_cast(coherence)); + int32x4_t tmp3 = + llvm_amdgcn_raw_buffer_load_i32x4(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset + 12 * sizeof(int32_t), + static_cast(coherence)); + + thread_buffer tmp; + + tmp.template get_as()(number<0>{}) = tmp0; + tmp.template get_as()(number<1>{}) = tmp1; + tmp.template get_as()(number<2>{}) = tmp2; + tmp.template get_as()(number<3>{}) = tmp3; + + return bit_cast(tmp); + } +} + +#ifndef BUFFER_LOAD_USE_INLINEASM +#define BUFFER_LOAD_USE_INLINEASM 0 +#endif + +template +CK_TILE_DEVICE thread_buffer amd_buffer_load_impl(int32x4_t src_wave_buffer_resource, + index_t src_thread_addr_offset, + index_t src_wave_addr_offset) +{ + static_assert( + (std::is_same::value && (N == 1 || N == 2 || N == 4 || N == 8)) || + (std::is_same::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) || + (std::is_same::value && (N == 1 || N == 2 || N == 4 || N == 8)) || + (std::is_same::value && (N == 1 || N == 2 || N == 4 || N == 8)) || + (std::is_same::value && + (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) || + (std::is_same::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) || + (std::is_same::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) || + (std::is_same::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) || + (std::is_same::value && + (N == 1 || N == 2 || N == 4 || N == 8 || N == 16 || N == 32)), + "wrong! not implemented"); + + using rtn_type = thread_buffer; + + if constexpr(std::is_same::value) // fp32 + { + if constexpr(N == 1) + { + return bit_cast( + llvm_amdgcn_raw_buffer_load_fp32(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset, + static_cast(coherence))); + } + else if constexpr(N == 2) + { + return bit_cast( + llvm_amdgcn_raw_buffer_load_fp32x2(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset, + static_cast(coherence))); + } + else if constexpr(N == 4) + { + return bit_cast( + llvm_amdgcn_raw_buffer_load_fp32x4(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset, + static_cast(coherence))); + } + else if constexpr(N == 8) + { + thread_buffer tmp; + + tmp.template get_as()(number<0>{}) = + llvm_amdgcn_raw_buffer_load_fp32x4(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset, + static_cast(coherence)); + + tmp.template get_as()(number<1>{}) = + llvm_amdgcn_raw_buffer_load_fp32x4(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset + 4 * sizeof(float), + static_cast(coherence)); + + return tmp; + } + else if constexpr(N == 16) + { + thread_buffer tmp; + + tmp.template get_as()(number<0>{}) = + llvm_amdgcn_raw_buffer_load_fp32x4(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset, + static_cast(coherence)); + + tmp.template get_as()(number<1>{}) = + llvm_amdgcn_raw_buffer_load_fp32x4(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset + 4 * sizeof(float), + static_cast(coherence)); + + tmp.template get_as()(number<2>{}) = + llvm_amdgcn_raw_buffer_load_fp32x4(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset + 8 * sizeof(float), + static_cast(coherence)); + + tmp.template get_as()(number<3>{}) = + llvm_amdgcn_raw_buffer_load_fp32x4(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset + 12 * sizeof(float), + static_cast(coherence)); + + return tmp; + } + } + else if constexpr(std::is_same::value) // fp16 + { + if constexpr(N == 1) + { + return bit_cast( + llvm_amdgcn_raw_buffer_load_fp16(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset, + static_cast(coherence))); + } + else if constexpr(N == 2) + { + return bit_cast( + llvm_amdgcn_raw_buffer_load_fp16x2(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset, + static_cast(coherence))); + } + else if constexpr(N == 4) + { + return bit_cast( + llvm_amdgcn_raw_buffer_load_fp16x4(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset, + static_cast(coherence))); + } + else if constexpr(N == 8) + { + // use fp32 load to mimic fp16 load + fp32x4_t tmp = llvm_amdgcn_raw_buffer_load_fp32x4(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset, + static_cast(coherence)); + + return bit_cast(tmp); + } + } + else if constexpr(std::is_same::value) // bf16 + { + if constexpr(N == 1) + { + return bit_cast( + llvm_amdgcn_raw_buffer_load_i16(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset, + static_cast(coherence))); + } + else if constexpr(N == 2) + { + return bit_cast( + llvm_amdgcn_raw_buffer_load_i16x2(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset, + static_cast(coherence))); + } + else if constexpr(N == 4) + { + return bit_cast( + llvm_amdgcn_raw_buffer_load_i16x4(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset, + static_cast(coherence))); + } + else if constexpr(N == 8) + { + int32x4_t tmp = llvm_amdgcn_raw_buffer_load_i32x4(src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset, + static_cast(coherence)); + + return bit_cast(tmp); + } + } + else // other datatype + { + auto raw_data = amd_buffer_load_impl_with_bytes( + src_wave_buffer_resource, src_thread_addr_offset, src_wave_addr_offset); + + return bit_cast(raw_data); + } +} + +template +CK_TILE_DEVICE void amd_buffer_load_raw_impl(thread_buffer& dst, + int32x4_t src_wave_buffer_resource, + index_t src_thread_addr_offset, + index_t src_wave_addr_offset, + index_t src_linear_addr_offset, + index_t flag = 0, + bool_constant = {}) +{ + constexpr index_t bytes = sizeof(T) * N; + static_assert(bytes == 1 || bytes == 2 || bytes == 4 || bytes == 8 || bytes == 16, + "wrong! not supported by buffer_load instruction"); + + using type = thread_buffer; + if constexpr(oob_conditional_check) + { + buffer_load_if{}(dst, + src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset, + src_linear_addr_offset, + flag, + bool_constant{}); + } + else + { + buffer_load{}(dst, + src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset, + src_linear_addr_offset, + flag, + bool_constant{}); + } +} + +template +CK_TILE_DEVICE void amd_async_buffer_load_impl(T* smem, + int32x4_t src_wave_buffer_resource, + index_t src_thread_addr_offset, + index_t src_wave_addr_offset, + index_t src_immediate_addr_offset = 0, + bool_constant = {}) +{ + static_assert(sizeof(T) * N == 4, "wrong! not implemented vector size"); + + async_buffer_load_dword_v(smem, + src_wave_buffer_resource, + src_thread_addr_offset, + src_wave_addr_offset, + src_immediate_addr_offset, + 0, + bool_constant{}); +} + +template +CK_TILE_DEVICE void amd_async_buffer_load(CK_TILE_LDS_ADDR T* smem, + int32x4_t src_wave_buffer_resource, + index_t src_thread_addr_offset, + index_t src_wave_addr_offset, + index_t src_immediate_addr_offset = 0, + index_t flag = 0, + bool_constant = {}) +{ + static_assert(sizeof(T) * N == 4, "wrong! not implemented vector size"); + + if constexpr(oob_conditional_check) + { + index_t v_offset = flag ? v_offset : src_wave_buffer_resource[2]; + llvm_amdgcn_raw_buffer_load_lds(src_wave_buffer_resource, + smem, + sizeof(uint32_t), + v_offset, + src_wave_addr_offset, + src_immediate_addr_offset, + static_cast(coherence)); + } + else + { + llvm_amdgcn_raw_buffer_load_lds(src_wave_buffer_resource, + smem, + sizeof(uint32_t), + src_thread_addr_offset, + src_wave_addr_offset, + src_immediate_addr_offset, + static_cast(coherence)); + } +} + +template +CK_TILE_DEVICE void amd_buffer_store_impl_with_bytes(const thread_buffer src_thread_data, + int32x4_t dst_wave_buffer_resource, + index_t dst_thread_addr_offset, + index_t dst_wave_addr_offset) +{ + static_assert(N == 1 || N == 2 || N == 4 || N == 8 || N == 16 || N == 32 || N == 64, + "wrong! not implemented"); + + if constexpr(N == 1) + { + llvm_amdgcn_raw_buffer_store_i8(bit_cast(src_thread_data), + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + static_cast(coherence)); + } + else if constexpr(N == 2) + { + + llvm_amdgcn_raw_buffer_store_i16(bit_cast(src_thread_data), + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + static_cast(coherence)); + } + else if constexpr(N == 4) + { + llvm_amdgcn_raw_buffer_store_i32(bit_cast(src_thread_data), + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + static_cast(coherence)); + } + else if constexpr(N == 8) + { + llvm_amdgcn_raw_buffer_store_i32x2(bit_cast(src_thread_data), + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + static_cast(coherence)); + } + else if constexpr(N == 16) + { + llvm_amdgcn_raw_buffer_store_i32x4(bit_cast(src_thread_data), + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + static_cast(coherence)); + } + else if constexpr(N == 32) + { + llvm_amdgcn_raw_buffer_store_i32x4( + src_thread_data.template get_as()[number<0>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + static_cast(coherence)); + + llvm_amdgcn_raw_buffer_store_i32x4( + src_thread_data.template get_as()[number<1>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset + sizeof(int32_t) * 4, + static_cast(coherence)); + } + else if constexpr(N == 64) + { + llvm_amdgcn_raw_buffer_store_i32x4( + src_thread_data.template get_as()[number<0>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + static_cast(coherence)); + + llvm_amdgcn_raw_buffer_store_i32x4( + src_thread_data.template get_as()[number<1>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset + sizeof(int32_t) * 4, + static_cast(coherence)); + + llvm_amdgcn_raw_buffer_store_i32x4( + src_thread_data.template get_as()[number<2>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset + sizeof(int32_t) * 8, + static_cast(coherence)); + + llvm_amdgcn_raw_buffer_store_i32x4( + src_thread_data.template get_as()[number<3>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset + sizeof(int32_t) * 12, + static_cast(coherence)); + } +} + +template +CK_TILE_DEVICE void amd_buffer_store_impl(const thread_buffer src_thread_data, + int32x4_t dst_wave_buffer_resource, + index_t dst_thread_addr_offset, + index_t dst_wave_addr_offset) +{ + static_assert( + (std::is_same::value && (N == 1 || N == 2 || N == 4 || N == 8)) || + (std::is_same::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) || + (std::is_same::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) || + (std::is_same::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) || + (std::is_same::value && + (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) || + (std::is_same::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) || + (std::is_same::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) || + (std::is_same::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) || + (std::is_same::value && + (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)) || + (std::is_same::value && (N == 1 || N == 2 || N == 4 || N == 8 || N == 16)), + "wrong! not implemented"); + + if constexpr(std::is_same::value) // fp32 + { + if constexpr(N == 1) + { + llvm_amdgcn_raw_buffer_store_fp32(bit_cast(src_thread_data), + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + static_cast(coherence)); + } + else if constexpr(N == 2) + { + llvm_amdgcn_raw_buffer_store_fp32x2(bit_cast(src_thread_data), + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + static_cast(coherence)); + } + else if constexpr(N == 4) + { + llvm_amdgcn_raw_buffer_store_fp32x4(bit_cast(src_thread_data), + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + static_cast(coherence)); + } + else if constexpr(N == 8) + { + llvm_amdgcn_raw_buffer_store_fp32x4( + src_thread_data.template get_as()[number<0>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + static_cast(coherence)); + llvm_amdgcn_raw_buffer_store_fp32x4( + src_thread_data.template get_as()[number<1>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset + 4 * sizeof(float), + static_cast(coherence)); + } + } + else if constexpr(std::is_same::value) // fp16 + { + if constexpr(N == 1) + { + llvm_amdgcn_raw_buffer_store_fp16(bit_cast<_Float16>(src_thread_data), + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + static_cast(coherence)); + } + else if constexpr(N == 2) + { + llvm_amdgcn_raw_buffer_store_fp16x2(bit_cast(src_thread_data), + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + static_cast(coherence)); + } + else if constexpr(N == 4) + { + llvm_amdgcn_raw_buffer_store_fp16x4(bit_cast(src_thread_data), + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + static_cast(coherence)); + } + else if constexpr(N == 8) + { +#if 0 + thread_buffer tmp{src_thread_data}; + + llvm_amdgcn_raw_buffer_store_fp16x4(tmp.template get_as()[number<0>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + static_cast(coherence)); + + llvm_amdgcn_raw_buffer_store_fp16x4(tmp.template get_as()[number<1>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset + 4 * sizeof(fp16_t), + static_cast(coherence)); +#else + llvm_amdgcn_raw_buffer_store_fp32x4(bit_cast(src_thread_data), + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + static_cast(coherence)); +#endif + } + } + else if constexpr(std::is_same::value) // bf16 + { + if constexpr(N == 1) + { + llvm_amdgcn_raw_buffer_store_i16(bit_cast(src_thread_data), + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + static_cast(coherence)); + } + else if constexpr(N == 2) + { + llvm_amdgcn_raw_buffer_store_i16x2(bit_cast(src_thread_data), + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + static_cast(coherence)); + } + else if constexpr(N == 4) + { + llvm_amdgcn_raw_buffer_store_i16x4(bit_cast(src_thread_data), + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + static_cast(coherence)); + } + else if constexpr(N == 8) + { + llvm_amdgcn_raw_buffer_store_i16x4( + src_thread_data.template get_as()[number<0>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + static_cast(coherence)); + + llvm_amdgcn_raw_buffer_store_i16x4( + src_thread_data.template get_as()[number<1>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset + 4 * sizeof(bf16_t), + static_cast(coherence)); + } + } + else if constexpr(std::is_same::value) + { + if constexpr(N == 1) + { + llvm_amdgcn_raw_buffer_store_ui16(bit_cast(src_thread_data), + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + static_cast(coherence)); + } + else if constexpr(N == 2) + { + llvm_amdgcn_raw_buffer_store_ui16x2(bit_cast(src_thread_data), + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + static_cast(coherence)); + } + else if constexpr(N == 4) + { + llvm_amdgcn_raw_buffer_store_ui16x4(bit_cast(src_thread_data), + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + static_cast(coherence)); + } + else if constexpr(N == 8) + { + llvm_amdgcn_raw_buffer_store_ui16x4( + src_thread_data.template get_as()[number<0>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + static_cast(coherence)); + + llvm_amdgcn_raw_buffer_store_ui16x4( + src_thread_data.template get_as()[number<1>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset + 4 * sizeof(uint16_t), + static_cast(coherence)); + } + } + else + { + using r_t = thread_buffer; + + amd_buffer_store_impl_with_bytes(bit_cast(src_thread_data), + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset); + } +} + +template +CK_TILE_DEVICE void amd_buffer_store_raw_impl(const thread_buffer& dst_thread_data, + int32x4_t dst_wave_buffer_resource, + index_t dst_thread_addr_offset, + index_t dst_wave_addr_offset, + index_t dst_linear_addr_offset, + index_t is_valid_element = 1) +{ + constexpr index_t bytes = sizeof(T) * N; + static_assert(bytes == 1 || bytes == 2 || bytes == 4 || bytes == 8 || bytes == 16, + "wrong! not supported by buffer_store instruction"); + + using type = thread_buffer; + if constexpr(oob_conditional_check) + { + buffer_store_if{}(dst_thread_data, + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + dst_linear_addr_offset, + is_valid_element); + } + else + { + buffer_store{}(dst_thread_data, + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + dst_linear_addr_offset); + } +} + +template +CK_TILE_DEVICE void amd_buffer_atomic_add_impl(const thread_buffer& src_thread_data, + int32x4_t dst_wave_buffer_resource, + index_t dst_thread_addr_offset, + index_t dst_wave_addr_offset) +{ + static_assert((std::is_same::value && (N == 1 || N == 2 || N == 4)) || + (std::is_same::value && (N == 2 || N == 4 || N == 8)) || + (std::is_same::value && (N == 1 || N == 2 || N == 4)), + "wrong! not implemented"); + + if constexpr(std::is_same::value) + { + if constexpr(N == 1) + { + llvm_amdgcn_raw_buffer_atomic_add_fp32(bit_cast(src_thread_data), + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + 0); + } + else if constexpr(N == 2) + { + llvm_amdgcn_raw_buffer_atomic_add_fp32( + src_thread_data.template get_as()[number<0>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + 0); + + llvm_amdgcn_raw_buffer_atomic_add_fp32( + src_thread_data.template get_as()[number<1>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset + sizeof(float), + 0); + } + else if constexpr(N == 4) + { + llvm_amdgcn_raw_buffer_atomic_add_fp32( + src_thread_data.template get_as()[number<0>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + 0); + + llvm_amdgcn_raw_buffer_atomic_add_fp32( + src_thread_data.template get_as()[number<1>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset + sizeof(float), + 0); + + llvm_amdgcn_raw_buffer_atomic_add_fp32( + src_thread_data.template get_as()[number<2>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset + 2 * sizeof(float), + 0); + + llvm_amdgcn_raw_buffer_atomic_add_fp32( + src_thread_data.template get_as()[number<3>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset + 3 * sizeof(float), + 0); + } + } + else if constexpr(std::is_same::value) + { + if constexpr(N == 2) + { + llvm_amdgcn_raw_buffer_atomic_add_fp16x2(bit_cast(src_thread_data), + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + 0); + } + else if constexpr(N == 4) + { + static_for<0, 2, 1>{}([&](auto i) { + llvm_amdgcn_raw_buffer_atomic_add_fp16x2( + src_thread_data.template get_as()[i], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset + i * sizeof(fp16x2_t), + 0); + }); + } + else if constexpr(N == 8) + { + static_for<0, 4, 1>{}([&](auto i) { + llvm_amdgcn_raw_buffer_atomic_add_fp16x2( + src_thread_data.template get_as()[i], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset + i * sizeof(fp16x2_t), + 0); + }); + } + } + else if constexpr(std::is_same::value) + { + if constexpr(N == 1) + { + llvm_amdgcn_raw_buffer_atomic_add_i32(bit_cast(src_thread_data), + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + 0); + } + else if constexpr(N == 2) + { + llvm_amdgcn_raw_buffer_atomic_add_i32( + src_thread_data.template get_as()[number<0>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + 0); + + llvm_amdgcn_raw_buffer_atomic_add_i32( + src_thread_data.template get_as()[number<1>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset + sizeof(int32_t), + 0); + } + else if constexpr(N == 4) + { + llvm_amdgcn_raw_buffer_atomic_add_i32( + src_thread_data.template get_as()[number<0>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + 0); + + llvm_amdgcn_raw_buffer_atomic_add_i32( + src_thread_data.template get_as()[number<1>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset + sizeof(int32_t), + 0); + + llvm_amdgcn_raw_buffer_atomic_add_i32( + src_thread_data.template get_as()[number<2>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset + 2 * sizeof(int32_t), + 0); + + llvm_amdgcn_raw_buffer_atomic_add_i32( + src_thread_data.template get_as()[number<3>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset + 3 * sizeof(int32_t), + 0); + } + } +} + +template +CK_TILE_DEVICE void amd_buffer_atomic_max_impl(const thread_buffer src_thread_data, + int32x4_t dst_wave_buffer_resource, + index_t dst_thread_addr_offset, + index_t dst_wave_addr_offset) +{ + static_assert((std::is_same::value && (N == 1 || N == 2 || N == 4)), + "wrong! not implemented"); + if constexpr(std::is_same::value) + { + if constexpr(N == 1) + { + llvm_amdgcn_raw_buffer_atomic_max_fp64(bit_cast(src_thread_data), + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + 0); + } + else if constexpr(N == 2) + { + llvm_amdgcn_raw_buffer_atomic_max_fp64( + src_thread_data.template get_as()[number<0>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + 0); + + llvm_amdgcn_raw_buffer_atomic_max_fp64( + src_thread_data.template get_as()[number<1>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset + sizeof(double), + 0); + } + else if constexpr(N == 4) + { + llvm_amdgcn_raw_buffer_atomic_max_fp64( + src_thread_data.template get_as()[number<0>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset, + 0); + + llvm_amdgcn_raw_buffer_atomic_max_fp64( + src_thread_data.template get_as()[number<1>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset + sizeof(double), + 0); + + llvm_amdgcn_raw_buffer_atomic_max_fp64( + src_thread_data.template get_as()[number<2>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset + 2 * sizeof(double), + 0); + + llvm_amdgcn_raw_buffer_atomic_max_fp64( + src_thread_data.template get_as()[number<3>{}], + dst_wave_buffer_resource, + dst_thread_addr_offset, + dst_wave_addr_offset + 3 * sizeof(double), + 0); + } + } +} + +// buffer_load requires: +// 1) p_src_wave must point to global memory space +// 2) p_src_wave must be a wavewise pointer. +// It is user's responsibility to make sure that is true. +// oob_conditional_check : dynamic check if out-of-bound +template +CK_TILE_DEVICE thread_buffer +amd_buffer_load_invalid_element_return_zero(const T* p_src_wave, + index_t src_thread_element_offset, + bool src_thread_element_valid, + index_t src_element_space_size) +{ + const int32x4_t src_wave_buffer_resource = + make_wave_buffer_resource(p_src_wave, src_element_space_size * sizeof(T)); + + index_t src_thread_addr_offset = src_thread_element_offset * sizeof(T); + +#if CK_TILE_EXPERIMENTAL_USE_BUFFER_LOAD_OOB_CHECK_OFFSET_TRICK + uint32_t src_addr_shift = [&]() { + if constexpr(oob_conditional_check) + return src_thread_element_valid ? 0 : 0x80000000; + else + return 0; + }(); + return amd_buffer_load_impl( + src_wave_buffer_resource, src_addr_shift + src_thread_addr_offset, 0); +#else + thread_buffer tmp = + amd_buffer_load_impl(src_wave_buffer_resource, src_thread_addr_offset, 0); + if constexpr(oob_conditional_check) + return src_thread_element_valid ? tmp : thread_buffer{numeric::zero()}; + else + return tmp; +#endif +} + +// buffer_load requires: +// 1) p_src_wave must point to global memory space +// 2) p_src_wave must be a wavewise pointer. +// It is user's responsibility to make sure that is true. +template +CK_TILE_DEVICE thread_buffer +amd_buffer_load_invalid_element_return_customized_value(const T* p_src_wave, + index_t src_thread_element_offset, + bool src_thread_element_valid, + index_t src_element_space_size, + T customized_value) +{ + const int32x4_t src_wave_buffer_resource = + make_wave_buffer_resource(p_src_wave, src_element_space_size * sizeof(T)); + + index_t src_thread_addr_offset = src_thread_element_offset * sizeof(T); + + thread_buffer tmp = + amd_buffer_load_impl(src_wave_buffer_resource, src_thread_addr_offset, 0); + + if constexpr(oob_conditional_check) + return src_thread_element_valid ? tmp : thread_buffer{customized_value}; + else + return tmp; +} + +template +CK_TILE_DEVICE void amd_buffer_load_raw(thread_buffer& dst, + const T* p_src_wave, + index_t src_thread_element_offset, + index_t src_linear_element_offset, + index_t src_element_space_size, + index_t is_valid_element = 0, + bool_constant = {}) +{ + const int32x4_t src_wave_buffer_resource = + make_wave_buffer_resource(p_src_wave, src_element_space_size * sizeof(T)); + + index_t src_thread_addr_offset = src_thread_element_offset * sizeof(T); + index_t src_linear_addr_offset = src_linear_element_offset * sizeof(T); + + amd_buffer_load_raw_impl( + dst, + src_wave_buffer_resource, + src_thread_addr_offset, + 0, + src_linear_addr_offset, + is_valid_element, + bool_constant{}); +} + +// This version support buffer resource as input arg +template +CK_TILE_DEVICE void amd_buffer_load_raw(thread_buffer& dst, + const int32x4_t src_wave_buffer_resource, + index_t src_thread_element_offset, + index_t src_linear_element_offset, + index_t is_valid_element = 0, + bool_constant = {}) +{ + index_t src_thread_addr_offset = src_thread_element_offset * sizeof(T); + index_t src_linear_addr_offset = src_linear_element_offset * sizeof(T); + + amd_buffer_load_raw_impl( + dst, + src_wave_buffer_resource, + src_thread_addr_offset, + 0, + src_linear_addr_offset, + is_valid_element, + bool_constant{}); +} + +// unfortunately async copy can not make sure invalid data is zero inside LDS +// ... unless people manually write zero to LDS at the proper address. +// so not support invalid_element check for now. +// buffer_load OOB still working. +template +CK_TILE_DEVICE void amd_async_buffer_load_with_oob_raw(T* smem, + const T* p_src_wave, + index_t src_thread_element_offset, + index_t src_linear_element_offset, + index_t src_element_space_size, + bool_constant = {}) +{ + const int32x4_t src_wave_buffer_resource = + make_wave_buffer_resource(p_src_wave, src_element_space_size * sizeof(T)); + + index_t src_thread_addr_offset = src_thread_element_offset * sizeof(T); + index_t src_linear_addr_offset = src_linear_element_offset * sizeof(T); + + amd_async_buffer_load_impl(smem, + src_wave_buffer_resource, + src_thread_addr_offset, + 0, + src_linear_addr_offset, + bool_constant{}); +} + +// This version support buffer resource as input arg +template +CK_TILE_DEVICE void amd_async_buffer_load_with_oob_raw(T* smem, + const int32x4_t src_wave_buffer_resource, + index_t src_thread_element_offset, + index_t src_linear_element_offset, + bool_constant = {}) +{ + index_t src_thread_addr_offset = src_thread_element_offset * sizeof(T); + index_t src_linear_addr_offset = src_linear_element_offset * sizeof(T); + + amd_async_buffer_load_impl(smem, + src_wave_buffer_resource, + src_thread_addr_offset, + 0, + src_linear_addr_offset, + bool_constant{}); +} + +// This version support buffer resource as input arg +template +CK_TILE_DEVICE void amd_async_buffer_load_with_oob(CK_TILE_LDS_ADDR T* smem, + const int32x4_t src_wave_buffer_resource, + index_t src_thread_element_offset, + index_t src_linear_element_offset, + bool is_valid_element, + bool_constant = {}) +{ + index_t src_thread_addr_offset = src_thread_element_offset * sizeof(T); + index_t src_linear_addr_offset = src_linear_element_offset * sizeof(T); + + amd_async_buffer_load(smem, + src_wave_buffer_resource, + src_thread_addr_offset, + 0, + src_linear_addr_offset, + is_valid_element, + bool_constant{}); +} + +// buffer_store requires: +// 1) p_dst_wave must point to global memory +// 2) p_dst_wave must be a wavewise pointer. +// It is user's responsibility to make sure that is true. +template +CK_TILE_DEVICE void amd_buffer_store(const thread_buffer& src_thread_data, + T* p_dst_wave, + const index_t dst_thread_element_offset, + const bool dst_thread_element_valid, + const index_t dst_element_space_size) +{ + const int32x4_t dst_wave_buffer_resource = + make_wave_buffer_resource(p_dst_wave, dst_element_space_size * sizeof(T)); + + index_t dst_thread_addr_offset = dst_thread_element_offset * sizeof(T); + +#if CK_TILE_EXPERIMENTAL_USE_BUFFER_STORE_OOB_CHECK_OFFSET_TRICK + uint32_t dst_addr_shift = [&]() { + if constexpr(oob_conditional_check) + return dst_thread_element_valid ? 0 : 0x80000000; + else + return 0; + }(); + amd_buffer_store_impl( + src_thread_data, dst_wave_buffer_resource, dst_addr_shift + dst_thread_addr_offset, 0); +#else + if constexpr(oob_conditional_check) + { + if(dst_thread_element_valid) + { + amd_buffer_store_impl( + src_thread_data, dst_wave_buffer_resource, dst_thread_addr_offset, 0); + } + } + else + { + amd_buffer_store_impl( + src_thread_data, dst_wave_buffer_resource, dst_thread_addr_offset, 0); + } +#endif +} + +template +CK_TILE_DEVICE void amd_buffer_store_raw(const thread_buffer& src_thread_data, + T* p_dst_wave, + const index_t dst_thread_element_offset, + const index_t dst_linear_element_offset, + const bool dst_thread_element_valid, + const index_t dst_element_space_size) +{ + const int32x4_t dst_wave_buffer_resource = + make_wave_buffer_resource(p_dst_wave, dst_element_space_size * sizeof(T)); + + index_t dst_thread_addr_offset = dst_thread_element_offset * sizeof(T); + index_t dst_linear_addr_offset = dst_linear_element_offset * sizeof(T); + + amd_buffer_store_raw_impl(src_thread_data, + dst_wave_buffer_resource, + dst_thread_addr_offset, + 0, + dst_linear_addr_offset, + dst_thread_element_valid); +} + +// buffer_atomic_add requires: +// 1) p_dst_wave must point to global memory +// 2) p_dst_wave must be a wavewise pointer. +// It is user's responsibility to make sure that is true. +template +CK_TILE_DEVICE void amd_buffer_atomic_add(const thread_buffer& src_thread_data, + T* p_dst_wave, + const index_t dst_thread_element_offset, + const bool dst_thread_element_valid, + const index_t dst_element_space_size) +{ + const int32x4_t dst_wave_buffer_resource = + make_wave_buffer_resource(p_dst_wave, dst_element_space_size * sizeof(T)); + + index_t dst_thread_addr_offset = dst_thread_element_offset * sizeof(T); + +#if CK_TILE_EXPERIMENTAL_USE_BUFFER_ATOMIC_ADD_OOB_CHECK_OFFSET_TRICK + uint32_t dst_addr_shift = dst_thread_element_valid ? 0 : 0x80000000; + + amd_buffer_atomic_add_impl( + src_thread_data, dst_wave_buffer_resource, dst_addr_shift + dst_thread_addr_offset, 0); +#else + if(dst_thread_element_valid) + { + amd_buffer_atomic_add_impl( + src_thread_data, dst_wave_buffer_resource, dst_thread_addr_offset, 0); + } +#endif +} + +template +CK_TILE_DEVICE void amd_buffer_atomic_add_raw(const thread_buffer& src_thread_data, + T* p_dst_wave, + const index_t dst_thread_element_offset, + const index_t dst_linear_element_offset, + const bool dst_thread_element_valid, + const index_t dst_element_space_size, + bool_constant = {}) +{ + const int32x4_t dst_wave_buffer_resource = + make_wave_buffer_resource(p_dst_wave, dst_element_space_size * sizeof(T)); + + index_t dst_thread_addr_offset = dst_thread_element_offset * sizeof(T); + index_t dst_linear_addr_offset = dst_linear_element_offset * sizeof(T); + + if constexpr(oob_conditional_check) + { + buffer_atomic_add_if{}(src_thread_data, + dst_wave_buffer_resource, + dst_thread_addr_offset, + 0, + dst_linear_addr_offset, + dst_thread_element_valid); + } + else + { + buffer_atomic_add{}(src_thread_data, + dst_wave_buffer_resource, + dst_thread_addr_offset, + 0, + dst_linear_addr_offset, + 1); + } +} + +// buffer_atomic_max requires: +// 1) p_dst_wave must point to global memory +// 2) p_dst_wave must be a wavewise pointer. +// It is user's responsibility to make sure that is true. +template +CK_TILE_DEVICE void amd_buffer_atomic_max(const thread_buffer& src_thread_data, + T* p_dst_wave, + const index_t dst_thread_element_offset, + const bool dst_thread_element_valid, + const index_t dst_element_space_size) +{ + const int32x4_t dst_wave_buffer_resource = + make_wave_buffer_resource(p_dst_wave, dst_element_space_size * sizeof(T)); + + index_t dst_thread_addr_offset = dst_thread_element_offset * sizeof(T); + +#if CK_TILE_EXPERIMENTAL_USE_BUFFER_ATOMIC_MAX_OOB_CHECK_OFFSET_TRICK + uint32_t dst_addr_shift = dst_thread_element_valid ? 0 : 0x80000000; + + amd_buffer_atomic_max_impl( + src_thread_data, dst_wave_buffer_resource, dst_addr_shift + dst_thread_addr_offset, 0); +#else + if(dst_thread_element_valid) + { + amd_buffer_atomic_max_impl( + src_thread_data, dst_wave_buffer_resource, dst_thread_addr_offset, 0); + } +#endif +} + +template +CK_TILE_DEVICE void amd_direct_load_global_to_lds(const T* global_base_ptr, + const index_t global_offset, + T* lds_base_ptr, + const index_t lds_offset, + const bool is_valid, + const index_t src_element_space_size) +{ + // Direct loads require that each thread reads and writes exactly a single DWORD. + constexpr auto dword_bytes = 4; + constexpr auto bytes_per_thread = sizeof(T) * NumElemsPerThread; + static_assert(bytes_per_thread == dword_bytes); + + const uint32_t* global_ptr = + reinterpret_cast(reinterpret_cast(global_base_ptr)); + const int32x4_t src_resource = + make_wave_buffer_resource(global_ptr, src_element_space_size * sizeof(T)); + const index_t global_offset_bytes = is_valid ? global_offset * sizeof(T) : 0x80000000; + +#if CK_TILE_USE_AMD_LDS_DIRECT_LOAD_INLINE_ASM + T* lds_ptr = lds_base_ptr + lds_offset; + auto const lds_ptr_sgpr = + __builtin_amdgcn_readfirstlane((reinterpret_cast(lds_ptr))); + asm volatile("s_mov_b32 m0, %0; \n\t" + "buffer_load_dword %1, %2, 0 offen lds;\n\t" ::"s"(lds_ptr_sgpr), + "v"(global_offset_bytes), + "s"(src_resource) + : "memory"); +#else + // LDS pointer must be attributed with the LDS address space. + __attribute__((address_space(3))) uint32_t* lds_ptr = + reinterpret_cast<__attribute__((address_space(3))) uint32_t*>( + reinterpret_cast(lds_base_ptr + lds_offset)); + + llvm_amdgcn_raw_buffer_load_lds( + src_resource, lds_ptr, sizeof(uint32_t), global_offset_bytes, 0, 0, 0); +#endif +} + +} // namespace ck_tile + +#endif // CK_TILE_USE_BUFFER_ADDRESSING_BUILTIN diff --git a/include/ck_tile/core/tensor/buffer_view.hpp b/include/ck_tile/core/tensor/buffer_view.hpp index c2a093f1ab..bdcfbdd920 100644 --- a/include/ck_tile/core/tensor/buffer_view.hpp +++ b/include/ck_tile/core/tensor/buffer_view.hpp @@ -5,7 +5,11 @@ #include "ck_tile/core/config.hpp" #include "ck_tile/core/arch/arch.hpp" +#if __clang_major__ == 20 +#include "ck_tile/core/arch/amd_buffer_addressing_builtins.hpp" +#else #include "ck_tile/core/arch/amd_buffer_addressing.hpp" +#endif #include "ck_tile/core/arch/generic_memory_space_atomic.hpp" #include "ck_tile/core/container/array.hpp" #include "ck_tile/core/numeric/integer.hpp" From 8146e471f14ad61b623c8a6410b9eb13787df91f Mon Sep 17 00:00:00 2001 From: Illia Silin <98187287+illsilin@users.noreply.github.com> Date: Fri, 23 May 2025 14:58:25 -0700 Subject: [PATCH 5/5] fix the buffer intrinsic names for clang >=20 (#2228) --- .../ck_tile/18_flatmm/run_flatmm_example.inc | 74 +++--- .../amd_buffer_addressing_builtins.hpp | 20 +- .../arch/amd_buffer_addressing_builtins.hpp | 218 +++++++++--------- 3 files changed, 160 insertions(+), 152 deletions(-) diff --git a/example/ck_tile/18_flatmm/run_flatmm_example.inc b/example/ck_tile/18_flatmm/run_flatmm_example.inc index 15a9df2c0c..c191fff7d0 100644 --- a/example/ck_tile/18_flatmm/run_flatmm_example.inc +++ b/example/ck_tile/18_flatmm/run_flatmm_example.inc @@ -4,14 +4,22 @@ #include template -constexpr const char* DataTypeToString() { - if constexpr (std::is_same_v) { +constexpr const char* DataTypeToString() +{ + if constexpr(std::is_same_v) + { return "fp16"; - } else if constexpr (std::is_same_v) { + } + else if constexpr(std::is_same_v) + { return "fp8"; - } else if constexpr (std::is_same_v) { + } + else if constexpr(std::is_same_v) + { return "bf8"; - } else { + } + else + { return "unknown"; } } @@ -112,8 +120,9 @@ float invoke_flatmm(ck_tile::DeviceMem& a_dev_buf, args.stride_B = stride_B; args.stride_C = stride_C; - float ave_time = flatmm_calc( - args, ck_tile::stream_config{nullptr, true, 1, n_warmup, n_repeat}); + float ave_time = + flatmm_calc( + args, ck_tile::stream_config{nullptr, true, 1, n_warmup, n_repeat}); std::size_t flop = std::size_t(2) * M * N * K; std::size_t num_byte = @@ -121,18 +130,15 @@ float invoke_flatmm(ck_tile::DeviceMem& a_dev_buf, float tflops = static_cast(flop) / 1.E9 / ave_time; float gb_per_sec = num_byte / 1.E6 / ave_time; - std::cout << "Run Flatmm kernel with DataType = " << DataTypeToString() << " M =" << M << " N =" << N << " K =" << K - << " StrideA =" << stride_A << " StrideB =" << stride_B << " StrideC =" << stride_C - << " : " << ave_time << " ms, " << tflops << " TFlops, " << gb_per_sec << " GB/s, " - << std::endl; + std::cout << "Run Flatmm kernel with DataType = " << DataTypeToString() + << " M =" << M << " N =" << N << " K =" << K << " StrideA =" << stride_A + << " StrideB =" << stride_B << " StrideC =" << stride_C << " : " << ave_time + << " ms, " << tflops << " TFlops, " << gb_per_sec << " GB/s, " << std::endl; return ave_time; } -template +template int run_flatmm_example_with_layouts(int argc, char* argv[], const ALayout a_layout = ALayout{}, @@ -147,7 +153,7 @@ int run_flatmm_example_with_layouts(int argc, using BDataType = typename GemmBasicTypeConfig::BDataType; using CDataType = typename GemmBasicTypeConfig::CDataType; using AccDataType = typename GemmBasicTypeConfig::AccDataType; - + ck_tile::index_t M = arg_parser.get_int("m"); ck_tile::index_t N = arg_parser.get_int("n"); ck_tile::index_t K = arg_parser.get_int("k"); @@ -182,7 +188,7 @@ int run_flatmm_example_with_layouts(int argc, c_rslt_host.SetZero(); // do pre-shuffle - std::string mfma = arg_parser.get_str("prec"); + std::string mfma = arg_parser.get_str("prec"); #if defined(USING_MFMA_16x16x32) && defined(ENABLE_FP8) ck_tile::index_t mfma_type = 1; #else @@ -193,18 +199,18 @@ int run_flatmm_example_with_layouts(int argc, b_shuffle_dev_buf.ToDevice(b_shuffle_host.data()); invoke_flatmm( - a_dev_buf, - b_shuffle_dev_buf, - c_dev_buf, - M, - N, - K, - stride_A, - stride_B, - stride_C, - kbatch, - n_warmup, - n_repeat); + a_dev_buf, + b_shuffle_dev_buf, + c_dev_buf, + M, + N, + K, + stride_A, + stride_B, + stride_C, + kbatch, + n_warmup, + n_repeat); c_dev_buf.FromDevice(c_rslt_host.data()); bool pass = true; @@ -219,8 +225,9 @@ int run_flatmm_example_with_layouts(int argc, a_host, b_origin_host, c_ref_host); const float max_accumulated_value = *std::max_element(c_ref_host.mData.begin(), c_ref_host.mData.end()); - const auto rtol_atol = calculate_rtol_atol(K, kbatch, max_accumulated_value); - pass = ck_tile::check_err(c_rslt_host, + const auto rtol_atol = calculate_rtol_atol( + K, kbatch, max_accumulated_value); + pass = ck_tile::check_err(c_rslt_host, c_ref_host, "Error: Incorrect results!", rtol_atol.at(ck_tile::number<0>{}), @@ -277,8 +284,9 @@ int run_flatmm_example_with_layouts(int argc, c_gpu_ref_dev_buf.FromDevice(c_gpu_ref_host.data()); const float max_accumulated_value = *std::max_element(c_gpu_ref_host.mData.begin(), c_gpu_ref_host.mData.end()); - const auto rtol_atol = calculate_rtol_atol(K, kbatch, max_accumulated_value); - pass = ck_tile::check_err(c_rslt_host, + const auto rtol_atol = calculate_rtol_atol( + K, kbatch, max_accumulated_value); + pass = ck_tile::check_err(c_rslt_host, c_gpu_ref_host, "Error: Incorrect results!", rtol_atol.at(ck_tile::number<0>{}), diff --git a/include/ck/utility/amd_buffer_addressing_builtins.hpp b/include/ck/utility/amd_buffer_addressing_builtins.hpp index 19869906dc..296c1d44d7 100644 --- a/include/ck/utility/amd_buffer_addressing_builtins.hpp +++ b/include/ck/utility/amd_buffer_addressing_builtins.hpp @@ -80,7 +80,7 @@ __device__ half2_t llvm_amdgcn_raw_buffer_atomic_add_fp16x2( int32x4_t rsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.fadd.v2f16.v4i32"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.fadd.v2f16"); // buffer atomic-add i32 __device__ int32_t llvm_amdgcn_raw_buffer_atomic_add_i32( @@ -88,7 +88,7 @@ __device__ int32_t llvm_amdgcn_raw_buffer_atomic_add_i32( int32x4_t rsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.add.i32.v4i32"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.add.i32"); // buffer atomic-add fp32 __device__ float llvm_amdgcn_raw_buffer_atomic_add_fp32( @@ -96,15 +96,15 @@ __device__ float llvm_amdgcn_raw_buffer_atomic_add_fp32( int32x4_t rsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.fadd.f32.v4i32"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.fadd.f32"); // buffer atomic-add fp32 -__device__ double llvm_amdgcn_raw_buffer_atomic_max_fp64( - double vdata, - int32x4_t rsrc, // dst_wave_buffer_resource - int voffset, // dst_thread_addr_offset - int soffset, // dst_wave_addr_offset - int glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.fmax.f64.v4i32"); +__device__ double +llvm_amdgcn_raw_buffer_atomic_max_fp64(double vdata, + int32x4_t rsrc, // dst_wave_buffer_resource + int voffset, // dst_thread_addr_offset + int soffset, // dst_wave_addr_offset + int glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.fmax.f64"); // memory coherency bit for buffer store/load instruction // check ISA manual for each GFX target @@ -827,7 +827,7 @@ llvm_amdgcn_raw_buffer_load_lds(int32x4_t rsrc, index_t voffset, index_t soffset, index_t offset, - index_t aux) __asm("llvm.amdgcn.raw.buffer.load.lds.v4i32"); + index_t aux) __asm("llvm.amdgcn.raw.buffer.load.lds"); #ifndef __HIPCC_RTC__ template diff --git a/include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp b/include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp index 0b9956cd01..53a344c7b0 100644 --- a/include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp +++ b/include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp @@ -881,95 +881,95 @@ CK_TILE_DEVICE_EXTERN int8_t llvm_amdgcn_raw_buffer_load_i8(int32x4_t srsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.i8.v4i32"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.i8"); CK_TILE_DEVICE_EXTERN int8x2_t llvm_amdgcn_raw_buffer_load_i8x2(int32x4_t srsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v2i8.v4i32"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v2i8"); CK_TILE_DEVICE_EXTERN int8x4_t llvm_amdgcn_raw_buffer_load_i8x4(int32x4_t srsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v4i8.v4i32"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v4i8"); // buffer load i16 CK_TILE_DEVICE_EXTERN int16_t llvm_amdgcn_raw_buffer_load_i16(int32x4_t srsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.i16.v4i32"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.i16"); CK_TILE_DEVICE_EXTERN int16x2_t llvm_amdgcn_raw_buffer_load_i16x2(int32x4_t srsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v2i16.v4i32"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v2i16"); CK_TILE_DEVICE_EXTERN int16x4_t llvm_amdgcn_raw_buffer_load_i16x4(int32x4_t srsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v4i16.v4i32"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v4i16"); // buffer load i32 CK_TILE_DEVICE_EXTERN int32_t llvm_amdgcn_raw_buffer_load_i32(int32x4_t srsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.i32.v4i32"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.i32"); CK_TILE_DEVICE_EXTERN int32x2_t llvm_amdgcn_raw_buffer_load_i32x2(int32x4_t srsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v2i32.v4i32"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v2i32"); CK_TILE_DEVICE_EXTERN int32x4_t llvm_amdgcn_raw_buffer_load_i32x4(int32x4_t srsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v4i32.v4i32"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v4i32"); // buffer load fp16 CK_TILE_DEVICE_EXTERN _Float16 llvm_amdgcn_raw_buffer_load_fp16(int32x4_t srsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.f16.v4i32"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.f16"); -CK_TILE_DEVICE_EXTERN fp16x2_t llvm_amdgcn_raw_buffer_load_fp16x2( - int32x4_t srsrc, - index_t voffset, - index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v2f16.v4i32"); +CK_TILE_DEVICE_EXTERN fp16x2_t +llvm_amdgcn_raw_buffer_load_fp16x2(int32x4_t srsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v2f16"); -CK_TILE_DEVICE_EXTERN fp16x4_t llvm_amdgcn_raw_buffer_load_fp16x4( - int32x4_t srsrc, - index_t voffset, - index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v4f16.v4i32"); +CK_TILE_DEVICE_EXTERN fp16x4_t +llvm_amdgcn_raw_buffer_load_fp16x4(int32x4_t srsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v4f16"); // buffer load fp32 CK_TILE_DEVICE_EXTERN float llvm_amdgcn_raw_buffer_load_fp32(int32x4_t srsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.f32.v4i32"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.f32"); -CK_TILE_DEVICE_EXTERN fp32x2_t llvm_amdgcn_raw_buffer_load_fp32x2( - int32x4_t srsrc, - index_t voffset, - index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v2f32.v4i32"); +CK_TILE_DEVICE_EXTERN fp32x2_t +llvm_amdgcn_raw_buffer_load_fp32x2(int32x4_t srsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v2f32"); -CK_TILE_DEVICE_EXTERN fp32x4_t llvm_amdgcn_raw_buffer_load_fp32x4( - int32x4_t srsrc, - index_t voffset, - index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v4f32.v4i32"); +CK_TILE_DEVICE_EXTERN fp32x4_t +llvm_amdgcn_raw_buffer_load_fp32x4(int32x4_t srsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.load.v4f32"); // buffer store i8 CK_TILE_DEVICE_EXTERN void @@ -977,21 +977,21 @@ llvm_amdgcn_raw_buffer_store_i8(int8_t vdata, int32x4_t rsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.i8.v4i32"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.i8"); CK_TILE_DEVICE_EXTERN void llvm_amdgcn_raw_buffer_store_i8x2(int8x2_t vdata, int32x4_t rsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2i8.v4i32"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2i8"); CK_TILE_DEVICE_EXTERN void llvm_amdgcn_raw_buffer_store_i8x4(int8x4_t vdata, int32x4_t rsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4i8.v4i32"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4i8"); // buffer store i16 CK_TILE_DEVICE_EXTERN void @@ -999,21 +999,21 @@ llvm_amdgcn_raw_buffer_store_i16(int16_t vdata, int32x4_t rsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.i16.v4i32"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.i16"); -CK_TILE_DEVICE_EXTERN void llvm_amdgcn_raw_buffer_store_i16x2( - int16x2_t vdata, - int32x4_t rsrc, - index_t voffset, - index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2i16.v4i32"); +CK_TILE_DEVICE_EXTERN void +llvm_amdgcn_raw_buffer_store_i16x2(int16x2_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2i16"); -CK_TILE_DEVICE_EXTERN void llvm_amdgcn_raw_buffer_store_i16x4( - int16x4_t vdata, - int32x4_t rsrc, - index_t voffset, - index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4i16.v4i32"); +CK_TILE_DEVICE_EXTERN void +llvm_amdgcn_raw_buffer_store_i16x4(int16x4_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4i16"); // buffer store i32 CK_TILE_DEVICE_EXTERN void @@ -1021,7 +1021,7 @@ llvm_amdgcn_raw_buffer_store_i32(int32_t vdata, int32x4_t rsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.i32.v4i32"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.i32"); // buffer store ui16 CK_TILE_DEVICE_EXTERN void @@ -1029,35 +1029,35 @@ llvm_amdgcn_raw_buffer_store_ui16(uint16_t vdata, int32x4_t rsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.i16.v4i32"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.i16"); -CK_TILE_DEVICE_EXTERN void llvm_amdgcn_raw_buffer_store_ui16x2( - uint16x2_t vdata, - int32x4_t rsrc, - index_t voffset, - index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2i16.v4i32"); +CK_TILE_DEVICE_EXTERN void +llvm_amdgcn_raw_buffer_store_ui16x2(uint16x2_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2i16"); -CK_TILE_DEVICE_EXTERN void llvm_amdgcn_raw_buffer_store_ui16x4( - uint16x4_t vdata, - int32x4_t rsrc, - index_t voffset, - index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4i16.v4i32"); +CK_TILE_DEVICE_EXTERN void +llvm_amdgcn_raw_buffer_store_ui16x4(uint16x4_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4i16"); -CK_TILE_DEVICE_EXTERN void llvm_amdgcn_raw_buffer_store_i32x2( - int32x2_t vdata, - int32x4_t rsrc, - index_t voffset, - index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2i32.v4i32"); +CK_TILE_DEVICE_EXTERN void +llvm_amdgcn_raw_buffer_store_i32x2(int32x2_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2i32"); -CK_TILE_DEVICE_EXTERN void llvm_amdgcn_raw_buffer_store_i32x4( - int32x4_t vdata, - int32x4_t rsrc, - index_t voffset, - index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4i32.v4i32"); +CK_TILE_DEVICE_EXTERN void +llvm_amdgcn_raw_buffer_store_i32x4(int32x4_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4i32"); // buffer store fp16 CK_TILE_DEVICE_EXTERN void @@ -1065,21 +1065,21 @@ llvm_amdgcn_raw_buffer_store_fp16(_Float16 vdata, int32x4_t rsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.f16.v4i32"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.f16"); -CK_TILE_DEVICE_EXTERN void llvm_amdgcn_raw_buffer_store_fp16x2( - fp16x2_t vdata, - int32x4_t rsrc, - index_t voffset, - index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2f16.v4i32"); +CK_TILE_DEVICE_EXTERN void +llvm_amdgcn_raw_buffer_store_fp16x2(fp16x2_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2f16"); -CK_TILE_DEVICE_EXTERN void llvm_amdgcn_raw_buffer_store_fp16x4( - fp16x4_t vdata, - int32x4_t rsrc, - index_t voffset, - index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4f16.v4i32"); +CK_TILE_DEVICE_EXTERN void +llvm_amdgcn_raw_buffer_store_fp16x4(fp16x4_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4f16"); // buffer store fp32 CK_TILE_DEVICE_EXTERN void @@ -1087,21 +1087,21 @@ llvm_amdgcn_raw_buffer_store_fp32(float vdata, int32x4_t rsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.f32.v4i32"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.f32"); -CK_TILE_DEVICE_EXTERN void llvm_amdgcn_raw_buffer_store_fp32x2( - fp32x2_t vdata, - int32x4_t rsrc, - index_t voffset, - index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2f32.v4i32"); +CK_TILE_DEVICE_EXTERN void +llvm_amdgcn_raw_buffer_store_fp32x2(fp32x2_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v2f32"); -CK_TILE_DEVICE_EXTERN void llvm_amdgcn_raw_buffer_store_fp32x4( - fp32x4_t vdata, - int32x4_t rsrc, - index_t voffset, - index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4f32.v4i32"); +CK_TILE_DEVICE_EXTERN void +llvm_amdgcn_raw_buffer_store_fp32x4(fp32x4_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4f32"); // buffer atomic-add fp16 CK_TILE_DEVICE_EXTERN fp16x2_t llvm_amdgcn_raw_buffer_atomic_add_fp16x2( @@ -1109,7 +1109,7 @@ CK_TILE_DEVICE_EXTERN fp16x2_t llvm_amdgcn_raw_buffer_atomic_add_fp16x2( int32x4_t rsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.fadd.v2f16.v4i32"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.fadd.v2f16"); // buffer atomic-add i32 CK_TILE_DEVICE_EXTERN int32_t llvm_amdgcn_raw_buffer_atomic_add_i32( @@ -1117,7 +1117,7 @@ CK_TILE_DEVICE_EXTERN int32_t llvm_amdgcn_raw_buffer_atomic_add_i32( int32x4_t rsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.add.i32.v4i32"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.add.i32"); // buffer atomic-add fp32 CK_TILE_DEVICE_EXTERN float llvm_amdgcn_raw_buffer_atomic_add_fp32( @@ -1125,15 +1125,15 @@ CK_TILE_DEVICE_EXTERN float llvm_amdgcn_raw_buffer_atomic_add_fp32( int32x4_t rsrc, index_t voffset, index_t soffset, - index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.fadd.f32.v4i32"); + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.fadd.f32"); // buffer atomic-max fp64 -CK_TILE_DEVICE_EXTERN double llvm_amdgcn_raw_buffer_atomic_max_fp64( - double vdata, - int32x4_t rsrc, // dst_wave_buffer_resource - int voffset, // dst_thread_addr_offset - int soffset, // dst_wave_addr_offset - int glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.fmax.f64.v4i32"); +CK_TILE_DEVICE_EXTERN double +llvm_amdgcn_raw_buffer_atomic_max_fp64(double vdata, + int32x4_t rsrc, // dst_wave_buffer_resource + int voffset, // dst_thread_addr_offset + int soffset, // dst_wave_addr_offset + int glc_slc) __asm("llvm.amdgcn.raw.buffer.atomic.fmax.f64"); // Direct loads from global to LDS. CK_TILE_DEVICE_EXTERN void @@ -1143,7 +1143,7 @@ llvm_amdgcn_raw_buffer_load_lds(int32x4_t rsrc, index_t voffset, index_t soffset, index_t offset, - index_t aux) __asm("llvm.amdgcn.raw.buffer.load.lds.v4i32"); + index_t aux) __asm("llvm.amdgcn.raw.buffer.load.lds"); template CK_TILE_DEVICE void async_buffer_load_dword_v(void* smem,