ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-12 17:26:00 +00:00
Code Issues Packages Projects Releases Wiki Activity

Rename enum

Browse Source 
Rename 'cood_transform_enum' to 'coord_transform_enum'
This commit is contained in:
Po-Yen, Chen
2024-03-13 03:15:04 -04:00
parent d962a0044b
commit 958218e9d0
3 changed files with 50 additions and 41 deletions
Download Patch File Download Diff File Expand all files Collapse all files

29
include/ck_tile/core/algorithm/coordinate_transform.hpp
View File

@@ -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 <typename LowLength>
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<LowLengths::size(), 1>
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 <typename LowLengths, typename RightShift>
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_; }

52
include/ck_tile/core/tensor/tensor_adaptor.hpp
View File

@@ -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<index_t>(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<index_t>(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<array<index_t, num_up_dim>>(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<array<index_t, num_low_dim>>(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<array<index_t, num_up_dim>>(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<array<index_t, num_up_dim>>(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<index_t>(0); \
\
return make_pass_through_transform(number<low_len>{}); \
} \
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<index_t>(0); \
\
@@ -876,7 +876,7 @@ CK_TILE_HOST_DEVICE constexpr auto chain_tensor_adaptors(const X& x, const Xs&..
return make_pad_transform( \
number<low_len>{}, number<left_pad>{}, number<right_pad>{}); \
} \
else if constexpr(name == cood_transform_enum::embed) \
else if constexpr(name == coord_transform_enum::embed) \
{ \
constexpr auto up_lens = \
meta_data.template get<array<index_t, num_up_dim>>(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<array<index_t, num_low_dim>>(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<array<index_t, num_up_dim>>(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<array<index_t, num_up_dim>>(0); \

10
include/ck_tile/core/tensor/tile_distribution.hpp
View File

@@ -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<kMaxMetaDataSize>;
using NumDim = index_t;
using Dims = array<index_t, kMaxNumDim>;
@@ -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<index_t, ndim_r_minor>(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<index_t, ndim_h_minor>(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<index_t, ndim_low>(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<index_t, ndim_y>(y_lengths)},
NumDim{1},
Dims{0},