diff --git a/include/ck_tile/core/algorithm/coordinate_transform.hpp b/include/ck_tile/core/algorithm/coordinate_transform.hpp index 7fa7eb9590..d0cd88ffbc 100644 --- a/include/ck_tile/core/algorithm/coordinate_transform.hpp +++ b/include/ck_tile/core/algorithm/coordinate_transform.hpp @@ -12,7 +12,7 @@ namespace ck_tile { -enum struct cood_transform_enum +enum struct coord_transform_enum { undefined, pass_through, @@ -30,7 +30,7 @@ struct base_transform { CK_TILE_HOST_DEVICE static constexpr auto get_type_enum() { - return cood_transform_enum::undefined; + return coord_transform_enum::undefined; } CK_TILE_HOST_DEVICE static constexpr index_t get_num_of_lower_dimension() { return NDimLow; } @@ -62,7 +62,7 @@ struct base_transform template struct pass_through : public base_transform<1, 1> { - static constexpr auto type_enum = cood_transform_enum::pass_through; + static constexpr auto type_enum = coord_transform_enum::pass_through; using LowerIndex = multi_index<1>; using UpperIndex = multi_index<1>; @@ -80,7 +80,7 @@ struct pass_through : public base_transform<1, 1> CK_TILE_HOST_DEVICE static constexpr auto get_type_enum() { - return cood_transform_enum::pass_through; + return coord_transform_enum::pass_through; } CK_TILE_HOST_DEVICE constexpr const auto& get_upper_lengths() const { return up_lengths_; } @@ -474,7 +474,10 @@ struct embed : public base_transform<1, UpLengths::size()> { } - CK_TILE_HOST_DEVICE static constexpr auto get_type_enum() { return cood_transform_enum::embed; } + CK_TILE_HOST_DEVICE static constexpr auto get_type_enum() + { + return coord_transform_enum::embed; + } CK_TILE_HOST_DEVICE constexpr const auto& get_upper_lengths() const { return up_lengths_; } @@ -602,7 +605,10 @@ struct merge_v2_magic_division : public base_transform static_assert(LowerIndex::size() == NDimLow, "wrong!"); } - CK_TILE_HOST_DEVICE static constexpr auto get_type_enum() { return cood_transform_enum::merge; } + CK_TILE_HOST_DEVICE static constexpr auto get_type_enum() + { + return coord_transform_enum::merge; + } CK_TILE_HOST_DEVICE constexpr const auto& get_upper_lengths() const { return up_lengths_; } @@ -870,7 +876,7 @@ struct unmerge : public base_transform<1, UpLengths::size()> CK_TILE_HOST_DEVICE static constexpr auto get_type_enum() { - return cood_transform_enum::unmerge; + return coord_transform_enum::unmerge; } CK_TILE_HOST_DEVICE constexpr const auto& get_upper_lengths() const { return up_lengths_; } @@ -1338,7 +1344,7 @@ struct modulo : public base_transform<1, 1> template struct xor_t : public base_transform<2, 2> { - static constexpr auto type_enum = cood_transform_enum::xor_t; + static constexpr auto type_enum = coord_transform_enum::xor_t; using LowerIndex = multi_index<2>; using UpperIndex = multi_index<2>; @@ -1356,7 +1362,10 @@ struct xor_t : public base_transform<2, 2> { } - CK_TILE_HOST_DEVICE static constexpr auto get_type_enum() { return cood_transform_enum::xor_t; } + CK_TILE_HOST_DEVICE static constexpr auto get_type_enum() + { + return coord_transform_enum::xor_t; + } CK_TILE_HOST_DEVICE constexpr const auto& get_upper_lengths() const { return up_lengths_; } @@ -1472,7 +1481,7 @@ struct offset : public base_transform<1, 1> CK_TILE_HOST_DEVICE static constexpr auto get_type_enum() { - return cood_transform_enum::offset; + return coord_transform_enum::offset; } CK_TILE_HOST_DEVICE constexpr const auto& get_upper_lengths() const { return up_lengths_; } diff --git a/include/ck_tile/core/tensor/tensor_adaptor.hpp b/include/ck_tile/core/tensor/tensor_adaptor.hpp index 8995dfd403..8a5c7218c6 100644 --- a/include/ck_tile/core/tensor/tensor_adaptor.hpp +++ b/include/ck_tile/core/tensor/tensor_adaptor.hpp @@ -695,7 +695,7 @@ CK_TILE_HOST_DEVICE constexpr auto chain_tensor_adaptors(const X& x, const Xs&.. // construct constexpr tensor_adaptor from constexpr encoding // encoded_tensor_adaptor are Tuple of following objects: // 1. encoded transforms (array of fixed size). Each encoded transform is a Tuple of following: -// 1.1 name (cood_transform_enum) +// 1.1 name (coord_transform_enum) // 1.2 meta data for constructor of the transform // 1.3 num of lower dimension (index_t) // 1.4 lower dimension Ids (array of fixed size) @@ -725,22 +725,22 @@ CK_TILE_HOST_DEVICE constexpr auto chain_tensor_adaptors(const X& x, const Xs&.. constexpr auto num_low_dim = encoded_transforms[i].template at<2>(); \ constexpr auto num_up_dim = encoded_transforms[i].template at<4>(); \ \ - STATIC_ASSERT(name == cood_transform_enum::pass_through || \ - name == cood_transform_enum::pad || \ - name == cood_transform_enum::embed || \ - name == cood_transform_enum::merge || \ - name == cood_transform_enum::unmerge || \ - name == cood_transform_enum::replicate, \ + STATIC_ASSERT(name == coord_transform_enum::pass_through || \ + name == coord_transform_enum::pad || \ + name == coord_transform_enum::embed || \ + name == coord_transform_enum::merge || \ + name == coord_transform_enum::unmerge || \ + name == coord_transform_enum::replicate, \ ""); \ \ - if constexpr(name == cood_transform_enum::pass_through) \ + if constexpr(name == coord_transform_enum::pass_through) \ { \ index_t pos = 0; \ auto low_len = meta_data.template pop(pos); \ \ return make_pass_through_transform(low_len); \ } \ - else if constexpr(name == cood_transform_enum::pad) \ + else if constexpr(name == coord_transform_enum::pad) \ { \ index_t pos = 0; \ auto low_len = meta_data.template pop(pos); \ @@ -749,7 +749,7 @@ CK_TILE_HOST_DEVICE constexpr auto chain_tensor_adaptors(const X& x, const Xs&.. \ return make_pad_transform(low_len, left_pad, right_pad); \ } \ - else if constexpr(name == cood_transform_enum::embed) \ + else if constexpr(name == coord_transform_enum::embed) \ { \ index_t pos = 0; \ auto up_lens = meta_data.template pop>(pos); \ @@ -758,21 +758,21 @@ CK_TILE_HOST_DEVICE constexpr auto chain_tensor_adaptors(const X& x, const Xs&.. \ return make_embed_transform(up_lens, coefficients); \ } \ - else if constexpr(name == cood_transform_enum::merge) \ + else if constexpr(name == coord_transform_enum::merge) \ { \ index_t pos = 0; \ auto low_lens = meta_data.template pop>(pos); \ \ return make_merge_transform(low_lens); \ } \ - else if constexpr(name == cood_transform_enum::unmerge) \ + else if constexpr(name == coord_transform_enum::unmerge) \ { \ index_t pos = 0; \ auto up_lens = meta_data.template pop>(pos); \ \ return make_unmerge_transform(up_lens); \ } \ - else if constexpr(name == cood_transform_enum::replicate) \ + else if constexpr(name == coord_transform_enum::replicate) \ { \ index_t pos = 0; \ auto up_lens = meta_data.template pop>(pos); \ @@ -819,7 +819,7 @@ CK_TILE_HOST_DEVICE constexpr auto chain_tensor_adaptors(const X& x, const Xs&.. // construct static tensor_adaptor from constexpr encoding // encoded_tensor_adaptor are Tuple of following objects: // 1. encoded transforms (array of fixed size). Each encoded transform is a Tuple of following: -// 1.1 name (cood_transform_enum) +// 1.1 name (coord_transform_enum) // 1.2 meta data for constructor of the transform // 1.3 num of lower dimension (index_t) // 1.4 lower dimension Ids (array of fixed size) @@ -849,21 +849,21 @@ CK_TILE_HOST_DEVICE constexpr auto chain_tensor_adaptors(const X& x, const Xs&.. constexpr auto num_low_dim = encoded_transforms[i].template at<2>(); \ constexpr auto num_up_dim = encoded_transforms[i].template at<4>(); \ \ - STATIC_ASSERT(name == cood_transform_enum::pass_through || \ - name == cood_transform_enum::pad || \ - name == cood_transform_enum::embed || \ - name == cood_transform_enum::merge || \ - name == cood_transform_enum::unmerge || \ - name == cood_transform_enum::replicate, \ + STATIC_ASSERT(name == coord_transform_enum::pass_through || \ + name == coord_transform_enum::pad || \ + name == coord_transform_enum::embed || \ + name == coord_transform_enum::merge || \ + name == coord_transform_enum::unmerge || \ + name == coord_transform_enum::replicate, \ ""); \ \ - if constexpr(name == cood_transform_enum::pass_through) \ + if constexpr(name == coord_transform_enum::pass_through) \ { \ constexpr index_t low_len = meta_data.template get(0); \ \ return make_pass_through_transform(number{}); \ } \ - else if constexpr(name == cood_transform_enum::pad) \ + else if constexpr(name == coord_transform_enum::pad) \ { \ constexpr index_t low_len = meta_data.template get(0); \ \ @@ -876,7 +876,7 @@ CK_TILE_HOST_DEVICE constexpr auto chain_tensor_adaptors(const X& x, const Xs&.. return make_pad_transform( \ number{}, number{}, number{}); \ } \ - else if constexpr(name == cood_transform_enum::embed) \ + else if constexpr(name == coord_transform_enum::embed) \ { \ constexpr auto up_lens = \ meta_data.template get>(0); \ @@ -887,21 +887,21 @@ CK_TILE_HOST_DEVICE constexpr auto chain_tensor_adaptors(const X& x, const Xs&.. return make_embed_transform(TO_TUPLE_OF_NUMBER(up_lens, num_up_dim), \ TO_TUPLE_OF_NUMBER(coefficients, num_up_dim)); \ } \ - else if constexpr(name == cood_transform_enum::merge) \ + else if constexpr(name == coord_transform_enum::merge) \ { \ constexpr auto low_lens = \ meta_data.template get>(0); \ \ return make_merge_transform(TO_TUPLE_OF_NUMBER(low_lens, num_low_dim)); \ } \ - else if constexpr(name == cood_transform_enum::unmerge) \ + else if constexpr(name == coord_transform_enum::unmerge) \ { \ constexpr auto up_lens = \ meta_data.template get>(0); \ \ return make_unmerge_transform(TO_TUPLE_OF_NUMBER(up_lens, num_up_dim)); \ } \ - else if constexpr(name == cood_transform_enum::replicate) \ + else if constexpr(name == coord_transform_enum::replicate) \ { \ constexpr auto up_lens = \ meta_data.template get>(0); \ diff --git a/include/ck_tile/core/tensor/tile_distribution.hpp b/include/ck_tile/core/tensor/tile_distribution.hpp index 4ee6ef72f5..9fee2fd5c6 100644 --- a/include/ck_tile/core/tensor/tile_distribution.hpp +++ b/include/ck_tile/core/tensor/tile_distribution.hpp @@ -246,7 +246,7 @@ CK_TILE_HOST_DEVICE constexpr auto constexpr index_t kMaxMetaDataSize = 128; constexpr index_t kMaxNumDim = 10; - using Name = cood_transform_enum; + using Name = coord_transform_enum; using MetaData = meta_data_buffer; using NumDim = index_t; using Dims = array; @@ -273,7 +273,7 @@ CK_TILE_HOST_DEVICE constexpr auto constexpr auto r_minor_lengths = RsLengths{}; trans(num_tran++) = { - cood_transform_enum::replicate, + coord_transform_enum::replicate, MetaData{to_array(r_minor_lengths)}, NumDim{0}, Dims{}, @@ -303,7 +303,7 @@ CK_TILE_HOST_DEVICE constexpr auto constexpr index_t ndim_h_minor = h_minor_lengths.size(); trans(num_tran++) = { - cood_transform_enum::unmerge, + coord_transform_enum::unmerge, MetaData{to_array(h_minor_lengths)}, NumDim{1}, Dims{idim_x}, @@ -348,7 +348,7 @@ CK_TILE_HOST_DEVICE constexpr auto low_lengths(i) = rh_major_minor_to_hidden_lengths[rh_major][rh_minor]; } - trans(num_tran++) = {cood_transform_enum::merge, + trans(num_tran++) = {coord_transform_enum::merge, MetaData{to_array(low_lengths)}, NumDim{ndim_low}, low_dims, @@ -394,7 +394,7 @@ CK_TILE_HOST_DEVICE constexpr auto d_length *= y_length; } - auto tran = make_tuple(cood_transform_enum::unmerge, + auto tran = make_tuple(coord_transform_enum::unmerge, MetaData{to_array(y_lengths)}, NumDim{1}, Dims{0},