/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-6.4.3/include/ck_tile/core/algorithm/coordinate_transform.hpp Source File#
coordinate_transform.hpp
Go to the documentation of this file.
137 calculate_upper_dimension_safe_vector_length_strides(const LowVectorLengths& low_vector_lengths,
325 calculate_upper_dimension_safe_vector_length_strides(const LowVectorLengths& low_vector_lengths,
425 calculate_upper_dimension_safe_vector_length_strides(const LowVectorLengths& low_vector_lengths,
455 // 3) Tuple of mixture of index_t and number, which is known partially at run-time and partially
563 // Implementation of "merge" transformation primitive that uses magic-number-division to do lowering
569 // dividend would be bit-wise interpreted as uint32_t and magic number division implementation for
691 calculate_upper_dimension_safe_vector_length_strides(const LowVectorLengths& low_vector_lengths,
720 // Implementation of "merge" transformation primitive that uses division and mod. It is supposed to
721 // be used for low_lengths that are known at compile time and are power of 2, otherwise performance
822 calculate_upper_dimension_safe_vector_length_strides(const LowVectorLengths& low_vector_lengths,
941 calculate_upper_dimension_safe_vector_length_strides(const LowVectorLengths& low_vector_lengths,
1359 CK_TILE_HOST_DEVICE constexpr xor_t(const LowLengths& low_lengths) : up_lengths_{low_lengths} {}
1612 //*******************************************************************************************************
1620 template <typename LowLength, typename LeftPad, typename RightPad, bool SkipIsValidCheck = false>
1747 make_indexing_transform_with_adaptor(const UpLength& up_lengths, const IndexingAdaptor& iadaptor)
__host__ constexpr __device__ auto unmerge(const Layout< Shape, UnrolledDesc > &layout, const NewLengths &new_lengths, [[maybe_unused]] const NewIdxs &new_indexes)
Unmerge selected dim in layout.
Definition: layout_utils.hpp:474
__host__ constexpr __device__ index_t gcd(index_t x, index_t y)
Definition: math.hpp:154
Definition: cluster_descriptor.hpp:13
constexpr CK_TILE_HOST_DEVICE auto make_insert_transform(const UpperIndex &up_idx)
Definition: coordinate_transform.hpp:1692
constexpr CK_TILE_HOST_DEVICE auto container_reduce(const Container &x, Reduce reduce, Init init, number< IBegin >=number< 0 >{}, number< IEnd >=number< Container::size()>{}, number< IStep >=number< 1 >{})
Definition: container_helper.hpp:198
constexpr CK_TILE_HOST_DEVICE auto make_left_pad_transform(const LowLength &low_length, const LeftPadLength &left_pad_, bool_constant< SkipIsValidCheck >=bool_constant< false >{})
Definition: coordinate_transform.hpp:1632
@ replicate
@ undefined
@ indexing
@ pass_through
constexpr CK_TILE_HOST_DEVICE auto make_right_pad_transform(const LowLength &low_length, const RightPadLength &right_pad_, bool_constant< SkipIsValidCheck >=bool_constant< false >{})
Definition: coordinate_transform.hpp:1641
__host__ __device__ multiplies() -> multiplies< void, void >
FIXME: create macro to replace 'host device' and nothing more.
constexpr CK_TILE_HOST_DEVICE auto make_indexing_transform_with_adaptor(const UpLength &up_lengths, const IndexingAdaptor &iadaptor)
Definition: coordinate_transform.hpp:1747
constexpr CK_TILE_HOST_DEVICE auto make_offset_transform(const LowLength &low_length, const OffsetLength &offset_length)
Definition: coordinate_transform.hpp:1725
is_static< T > is_known_at_compile_time
Definition: type_traits.hpp:93
constexpr CK_TILE_HOST_DEVICE auto make_slice_transform(const LowLength &low_length, const SliceBegin &slice_begin, const SliceEnd &slice_end)
Definition: coordinate_transform.hpp:1704
constexpr CK_TILE_HOST_DEVICE auto make_merge_transform(const LowLengths &low_lengths)
Definition: coordinate_transform.hpp:1672
constexpr CK_TILE_HOST_DEVICE auto make_pass_through_transform(const LowLength &low_length)
Definition: coordinate_transform.hpp:1615
constexpr CK_TILE_HOST_DEVICE auto make_pad_transform(const LowLength &low_length, const LeftPad &left_pad, const RightPad &right_pad, bool_constant< SkipIsValidCheck >=bool_constant< false >{})
Definition: coordinate_transform.hpp:1622
constexpr CK_TILE_HOST_DEVICE auto make_unmerge_transform(const UpLengths &up_lengths, bool_constant< Use24BitIntegerCalculation >=bool_constant< false >{})
Definition: coordinate_transform.hpp:1679
constexpr CK_TILE_HOST_DEVICE auto make_merge_transform_v3_division_mod(const LowLengths &low_lengths)
Definition: coordinate_transform.hpp:1666
constexpr CK_TILE_HOST_DEVICE auto make_modulo_transform(const Modulus &modulus, const UpLength &up_length)
Definition: coordinate_transform.hpp:1712
constexpr CK_TILE_HOST_DEVICE auto make_indexing_transform(const UpLength &up_lengths, const Indices &indices)
Definition: coordinate_transform.hpp:1737
constexpr CK_TILE_HOST_DEVICE auto generate_tuple(F &&f, number< N >)
Definition: tuple.hpp:400
constexpr CK_TILE_HOST_DEVICE auto make_tuple(Xs &&... xs)
Definition: tuple.hpp:337
constexpr CK_TILE_HOST_DEVICE auto make_xor_transform(const LowLengths &low_lengths)
Definition: coordinate_transform.hpp:1719
constexpr CK_TILE_HOST_DEVICE auto make_replicate_transform(const UpLengths &up_lengths)
Definition: coordinate_transform.hpp:1698
constexpr CK_TILE_HOST_DEVICE auto make_freeze_transform(const LowerIndex &low_idx)
Definition: coordinate_transform.hpp:1686
constexpr CK_TILE_HOST_DEVICE auto make_merge_transform_v2_magic_division(const LowLengths &low_lengths)
Definition: coordinate_transform.hpp:1659
constexpr CK_TILE_HOST_DEVICE auto make_embed_transform(const UpLengths &up_lengths, const Coefficients &coefficients)
Definition: coordinate_transform.hpp:1651
__host__ constexpr __device__ auto container_reverse_exclusive_scan(const Array< TData, NSize > &x, Reduce f, TData init)
Definition: container_helper.hpp:213
__host__ constexpr __device__ auto container_reduce(const Container &x, Reduce reduce, Init init, Number< IBegin >=Number< 0 >{}, Number< IEnd >=Number< Container::Size()>{}, Number< IStep >=Number< 1 >{})
Definition: container_helper.hpp:111
Definition: array.hpp:24
Definition: coordinate_transform.hpp:31
static constexpr CK_TILE_HOST_DEVICE auto get_type_enum()
Definition: coordinate_transform.hpp:32
static constexpr CK_TILE_HOST_DEVICE index_t get_num_of_upper_dimension()
Definition: coordinate_transform.hpp:39
static constexpr CK_TILE_HOST_DEVICE index_t get_num_of_lower_dimension()
Definition: coordinate_transform.hpp:37
static constexpr CK_TILE_HOST_DEVICE auto calculate_upper_dimension_safe_vector_length_strides(const LowVectorLengths &, const LowVectorStrides &)
Definition: coordinate_transform.hpp:46
Definition: integral_constant.hpp:13
Definition: coordinate_transform.hpp:461
constexpr CK_TILE_HOST_DEVICE embed()=default
constexpr CK_TILE_HOST_DEVICE const auto & get_upper_lengths() const
Definition: coordinate_transform.hpp:483
static constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_always_mapped_to_valid_lower_index()
Definition: coordinate_transform.hpp:518
constexpr CK_TILE_HOST_DEVICE void calculate_lower_index(LowIdx &idx_low, const UpIdx &idx_up) const
Definition: coordinate_transform.hpp:486
CK_TILE_HOST_DEVICE void update_lower_index(LowIdxDiff &idx_diff_low, const UpIdxDiff &idx_diff_up, LowIdx &idx_low, const UpIdx &) const
Definition: coordinate_transform.hpp:500
static constexpr CK_TILE_HOST_DEVICE auto get_type_enum()
Definition: coordinate_transform.hpp:478
static constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_mapped_to_valid_lower_index(const UpIdx &)
Definition: coordinate_transform.hpp:525
static constexpr CK_TILE_HOST_DEVICE bool is_known_at_compile_time()
Definition: coordinate_transform.hpp:530
constexpr CK_TILE_HOST_DEVICE embed(const UpLengths &up_lengths, const Coefficients &coefficients)
Definition: coordinate_transform.hpp:472
Definition: coordinate_transform.hpp:981
constexpr CK_TILE_HOST_DEVICE void calculate_lower_index(LowIdx &idx_low, const UpIdx &) const
Definition: coordinate_transform.hpp:991
static CK_TILE_HOST_DEVICE void update_lower_index(LowIdxDiff &idx_diff_low, const UpIdxDiff &, LowIdx &, const UpIdx &)
Definition: coordinate_transform.hpp:1001
static constexpr CK_TILE_HOST_DEVICE auto get_upper_lengths()
Definition: coordinate_transform.hpp:988
constexpr CK_TILE_HOST_DEVICE freeze()=default
constexpr CK_TILE_HOST_DEVICE freeze(const LowerIndex &low_idx)
Definition: coordinate_transform.hpp:986
static constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_always_mapped_to_valid_lower_index()
Definition: coordinate_transform.hpp:1010
static constexpr CK_TILE_HOST_DEVICE bool is_known_at_compile_time()
Definition: coordinate_transform.hpp:1022
CK_TILE_HOST_DEVICE void print() const
Definition: coordinate_transform.hpp:1027
static constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_mapped_to_valid_lower_index(const UpIdx &)
Definition: coordinate_transform.hpp:1017
Definition: type_traits.hpp:75
Definition: indexing_adaptor.hpp:20
Definition: coordinate_transform.hpp:1532
static constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_mapped_to_valid_lower_index(const UpIdx &)
Definition: coordinate_transform.hpp:1588
constexpr CK_TILE_HOST_DEVICE indexing()=default
constexpr CK_TILE_HOST_DEVICE void calculate_lower_index(LowIdx &idx_low, const UpIdx &idx_up) const
Definition: coordinate_transform.hpp:1558
static constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_always_mapped_to_valid_lower_index()
Definition: coordinate_transform.hpp:1581
decltype(make_tuple(UpLength{})) UpLengths
Definition: coordinate_transform.hpp:1538
CK_TILE_HOST_DEVICE void print() const
Definition: coordinate_transform.hpp:1599
constexpr CK_TILE_HOST_DEVICE const auto & get_upper_lengths() const
Definition: coordinate_transform.hpp:1555
CK_TILE_HOST_DEVICE void update_lower_index(LowIdxDiff &idx_diff_low, const UpIdxDiff &idx_diff_up, LowIdx &idx_low, const UpIdx &idx_up) const
Definition: coordinate_transform.hpp:1567
static constexpr CK_TILE_HOST_DEVICE bool is_known_at_compile_time()
Definition: coordinate_transform.hpp:1593
constexpr CK_TILE_HOST_DEVICE indexing(const UpLength &up_length, const IndexingAdaptor &iadaptor)
Definition: coordinate_transform.hpp:1544
static constexpr CK_TILE_HOST_DEVICE auto get_type_enum()
Definition: coordinate_transform.hpp:1550
Definition: coordinate_transform.hpp:1042
static CK_TILE_HOST_DEVICE void update_lower_index(LowIdxDiff &, const UpIdxDiff &, LowIdx &, const UpIdx &)
Definition: coordinate_transform.hpp:1069
constexpr CK_TILE_HOST_DEVICE insert(const UpperLength &up_length)
Definition: coordinate_transform.hpp:1049
CK_TILE_HOST_DEVICE void print() const
Definition: coordinate_transform.hpp:1096
decltype(make_tuple(UpperLength{})) UpLengths
Definition: coordinate_transform.hpp:1043
static constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_mapped_to_valid_lower_index(const UpIdx &)
Definition: coordinate_transform.hpp:1086
constexpr CK_TILE_HOST_DEVICE insert()=default
constexpr CK_TILE_HOST_DEVICE void calculate_lower_index(LowIdx &, const UpIdx &) const
Definition: coordinate_transform.hpp:1061
static constexpr CK_TILE_HOST_DEVICE index_t get_num_of_lower_dimension()
Definition: coordinate_transform.hpp:1054
constexpr CK_TILE_HOST_DEVICE auto get_upper_lengths() const
Definition: coordinate_transform.hpp:1058
static constexpr CK_TILE_HOST_DEVICE bool IsLinearTransform()
Definition: coordinate_transform.hpp:1076
static constexpr CK_TILE_HOST_DEVICE bool is_known_at_compile_time()
Definition: coordinate_transform.hpp:1091
static constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_always_mapped_to_valid_lower_index()
Definition: coordinate_transform.hpp:1079
static constexpr CK_TILE_HOST_DEVICE index_t get_num_of_upper_dimension()
Definition: coordinate_transform.hpp:1056
Definition: coordinate_transform.hpp:555
constexpr CK_TILE_HOST_DEVICE auto operator()(number< I > i) const
Definition: coordinate_transform.hpp:557
Definition: coordinate_transform.hpp:257
static constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_always_mapped_to_valid_lower_index()
Definition: coordinate_transform.hpp:304
static constexpr CK_TILE_HOST_DEVICE bool is_known_at_compile_time()
Definition: coordinate_transform.hpp:316
LeftPadLength left_pad_length_
Definition: coordinate_transform.hpp:264
constexpr CK_TILE_HOST_DEVICE left_pad()=default
decltype(make_tuple(LowLength{}+LeftPadLength{})) UpLengths
Definition: coordinate_transform.hpp:261
constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_mapped_to_valid_lower_index(const UpIdx &idx_up) const
Definition: coordinate_transform.hpp:311
constexpr CK_TILE_HOST_DEVICE left_pad(const LowLength &low_length, const LeftPadLength &left_pad_length)
Definition: coordinate_transform.hpp:268
static CK_TILE_HOST_DEVICE void update_lower_index(LowIdxDiff &idx_diff_low, const UpIdxDiff &idx_diff_up, LowIdx &idx_low, const UpIdx &)
Definition: coordinate_transform.hpp:287
CK_TILE_HOST_DEVICE void print() const
Definition: coordinate_transform.hpp:334
constexpr CK_TILE_HOST_DEVICE const auto & get_upper_lengths() const
Definition: coordinate_transform.hpp:274
constexpr CK_TILE_HOST_DEVICE void calculate_lower_index(LowIdx &idx_low, const UpIdx &idx_up) const
Definition: coordinate_transform.hpp:277
static constexpr CK_TILE_HOST_DEVICE auto calculate_upper_dimension_safe_vector_length_strides(const LowVectorLengths &low_vector_lengths, const LowVectorStrides &low_vector_strides)
Definition: coordinate_transform.hpp:325
static constexpr CK_TILE_HOST_DEVICE auto calculate_magic_numbers(uint32_t divisor)
Definition: magic_div.hpp:29
static constexpr CK_TILE_DEVICE uint32_t do_magic_division(uint32_t dividend, uint32_t multiplier, uint32_t shift)
Definition: magic_div.hpp:60
Definition: coordinate_transform.hpp:577
static constexpr CK_TILE_HOST_DEVICE auto calculate_upper_dimension_safe_vector_length_strides(const LowVectorLengths &low_vector_lengths, const LowVectorStrides &low_vector_strides)
Definition: coordinate_transform.hpp:691
CK_TILE_HOST_DEVICE void update_lower_index(LowIdxDiff &idx_diff_low, const UpIdxDiff &, LowIdx &idx_low, const UpIdx &idx_up_new) const
Definition: coordinate_transform.hpp:638
LowLengthsMagicDivisor low_lengths_magic_divisor_
Definition: coordinate_transform.hpp:591
static constexpr auto I1
Definition: coordinate_transform.hpp:595
static constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_always_mapped_to_valid_lower_index()
Definition: coordinate_transform.hpp:669
static constexpr index_t NDimLow
Definition: coordinate_transform.hpp:578
static constexpr CK_TILE_HOST_DEVICE bool is_known_at_compile_time()
Definition: coordinate_transform.hpp:674
constexpr CK_TILE_HOST_DEVICE const auto & get_upper_lengths() const
Definition: coordinate_transform.hpp:614
decltype(generate_tuple(lambda_merge_generate_MagicDivision_calculate_magic_divisor< LowLengths >{}, number< NDimLow >{})) LowLengthsMagicDivisor
Definition: coordinate_transform.hpp:588
LowLengths low_lengths_
Definition: coordinate_transform.hpp:590
static constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_mapped_to_valid_lower_index(const UpIdx &)
Definition: coordinate_transform.hpp:683
static constexpr auto I0
Definition: coordinate_transform.hpp:594
UpLengths up_lengths_
Definition: coordinate_transform.hpp:592
decltype(make_tuple(container_reduce(LowLengths{}, multiplies{}, number< 1 >{}))) UpLengths
Definition: coordinate_transform.hpp:584
constexpr CK_TILE_HOST_DEVICE void calculate_lower_index(LowIdx &idx_low, const UpIdx &idx_up) const
Definition: coordinate_transform.hpp:617
CK_TILE_HOST_DEVICE void print() const
Definition: coordinate_transform.hpp:703
constexpr CK_TILE_HOST_DEVICE merge_v2_magic_division()=default
static constexpr CK_TILE_HOST_DEVICE auto get_type_enum()
Definition: coordinate_transform.hpp:609
constexpr CK_TILE_HOST_DEVICE merge_v2_magic_division(const LowLengths &low_lengths)
Definition: coordinate_transform.hpp:599
Definition: coordinate_transform.hpp:725
decltype(make_tuple(container_reduce(LowLengths{}, multiplies{}, number< 1 >{}))) UpLengths
Definition: coordinate_transform.hpp:735
static constexpr CK_TILE_HOST_DEVICE auto calculate_upper_dimension_safe_vector_length_strides(const LowVectorLengths &low_vector_lengths, const LowVectorStrides &low_vector_strides)
Definition: coordinate_transform.hpp:822
decltype(container_reverse_exclusive_scan(LowLengths{}, multiplies{}, number< 1 >{})) LowLengthsScan
Definition: coordinate_transform.hpp:732
LowLengths low_lengths_
Definition: coordinate_transform.hpp:737
static constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_mapped_to_valid_lower_index(const UpIdx &)
Definition: coordinate_transform.hpp:814
UpLengths up_lengths_
Definition: coordinate_transform.hpp:739
constexpr CK_TILE_HOST_DEVICE const auto & get_upper_lengths() const
Definition: coordinate_transform.hpp:752
constexpr CK_TILE_HOST_DEVICE merge_v3_division_mod(const LowLengths &low_lengths)
Definition: coordinate_transform.hpp:743
LowLengthsScan low_lengths_scan_
Definition: coordinate_transform.hpp:738
static constexpr CK_TILE_HOST_DEVICE bool is_known_at_compile_time()
Definition: coordinate_transform.hpp:805
constexpr CK_TILE_HOST_DEVICE void calculate_lower_index(LowIdx &idx_low, const UpIdx &idx_up) const
Definition: coordinate_transform.hpp:755
CK_TILE_HOST_DEVICE void update_lower_index(LowIdxDiff &idx_diff_low, const UpIdxDiff &, LowIdx &idx_low, const UpIdx &idx_up_new) const
Definition: coordinate_transform.hpp:773
CK_TILE_HOST_DEVICE void print() const
Definition: coordinate_transform.hpp:834
constexpr CK_TILE_HOST_DEVICE merge_v3_division_mod()=default
static constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_always_mapped_to_valid_lower_index()
Definition: coordinate_transform.hpp:800
Definition: coordinate_transform.hpp:1270
static constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_mapped_to_valid_lower_index(const UpIdx &)
Definition: coordinate_transform.hpp:1322
decltype(make_tuple(UpLength{})) UpLengths
Definition: coordinate_transform.hpp:1273
constexpr CK_TILE_HOST_DEVICE modulo(const Modulus &modulus, const UpLength &up_length)
Definition: coordinate_transform.hpp:1280
CK_TILE_HOST_DEVICE void update_lower_index(LowIdxDiff &idx_diff_low, const UpIdxDiff &idx_diff_up, LowIdx &idx_low, const UpIdx &up_idx) const
Definition: coordinate_transform.hpp:1298
constexpr CK_TILE_HOST_DEVICE void calculate_lower_index(LowIdx &idx_low, const UpIdx &idx_up) const
Definition: coordinate_transform.hpp:1288
constexpr CK_TILE_HOST_DEVICE const auto & get_upper_lengths() const
Definition: coordinate_transform.hpp:1285
CK_TILE_HOST_DEVICE void print() const
Definition: coordinate_transform.hpp:1332
constexpr CK_TILE_HOST_DEVICE modulo()=default
static constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_always_mapped_to_valid_lower_index()
Definition: coordinate_transform.hpp:1315
static constexpr CK_TILE_HOST_DEVICE bool is_known_at_compile_time()
Definition: coordinate_transform.hpp:1327
Definition: math.hpp:98
Definition: coordinate_transform.hpp:1443
static CK_TILE_HOST_DEVICE void update_lower_index(LowIdxDiff &idx_diff_low, const UpIdxDiff &idx_diff_up, LowIdx &idx_low, const UpIdx &)
Definition: coordinate_transform.hpp:1478
static constexpr CK_TILE_HOST_DEVICE bool is_known_at_compile_time()
Definition: coordinate_transform.hpp:1507
decltype(make_tuple(LowLength{})) UpLengths
Definition: coordinate_transform.hpp:1447
static constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_always_mapped_to_valid_lower_index()
Definition: coordinate_transform.hpp:1495
constexpr CK_TILE_HOST_DEVICE const auto & get_upper_lengths() const
Definition: coordinate_transform.hpp:1465
CK_TILE_HOST_DEVICE void print() const
Definition: coordinate_transform.hpp:1513
constexpr CK_TILE_HOST_DEVICE offset()=default
static constexpr CK_TILE_HOST_DEVICE auto get_type_enum()
Definition: coordinate_transform.hpp:1460
constexpr CK_TILE_HOST_DEVICE offset(const LowLength &low_length, const OffsetLength &offset_length)
Definition: coordinate_transform.hpp:1454
constexpr CK_TILE_HOST_DEVICE void calculate_lower_index(LowIdx &idx_low, const UpIdx &idx_up) const
Definition: coordinate_transform.hpp:1468
constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_mapped_to_valid_lower_index(const UpIdx &) const
Definition: coordinate_transform.hpp:1502
Definition: coordinate_transform.hpp:161
constexpr CK_TILE_HOST_DEVICE pad(const LowLength &low_length, const LeftPadLength &left_pad_length, const RightPadLength &right_pad_length)
Definition: coordinate_transform.hpp:173
decltype(make_tuple(LowLength{}+LeftPadLength{}+RightPadLength{})) UpLengths
Definition: coordinate_transform.hpp:165
constexpr CK_TILE_HOST_DEVICE const auto & get_upper_lengths() const
Definition: coordinate_transform.hpp:182
LeftPadLength left_pad_length_
Definition: coordinate_transform.hpp:168
static constexpr CK_TILE_HOST_DEVICE bool is_known_at_compile_time()
Definition: coordinate_transform.hpp:226
constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_mapped_to_valid_lower_index(const UpIdx &idx_up) const
Definition: coordinate_transform.hpp:219
static CK_TILE_HOST_DEVICE void update_lower_index(LowIdxDiff &idx_diff_low, const UpIdxDiff &idx_diff_up, LowIdx &idx_low, const UpIdx &)
Definition: coordinate_transform.hpp:195
constexpr CK_TILE_HOST_DEVICE void calculate_lower_index(LowIdx &idx_low, const UpIdx &idx_up) const
Definition: coordinate_transform.hpp:185
RightPadLength right_pad_length_
Definition: coordinate_transform.hpp:169
static constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_always_mapped_to_valid_lower_index()
Definition: coordinate_transform.hpp:212
Definition: coordinate_transform.hpp:65
static constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_mapped_to_valid_lower_index(const UpIdx &)
Definition: coordinate_transform.hpp:124
constexpr CK_TILE_HOST_DEVICE pass_through(const LowLength &low_length)
Definition: coordinate_transform.hpp:77
static constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_always_mapped_to_valid_lower_index()
Definition: coordinate_transform.hpp:117
decltype(make_tuple(LowLength{})) UpLengths
Definition: coordinate_transform.hpp:71
constexpr CK_TILE_HOST_DEVICE pass_through()=default
static constexpr auto type_enum
Definition: coordinate_transform.hpp:66
static CK_TILE_HOST_DEVICE void update_lower_index(LowIdxDiff &idx_diff_low, const UpIdxDiff &idx_diff_up, LowIdx &idx_low, const UpIdx &)
Definition: coordinate_transform.hpp:100
static constexpr CK_TILE_HOST_DEVICE void calculate_lower_index(LowIdx &idx_low, const UpIdx &idx_up)
Definition: coordinate_transform.hpp:90
CK_TILE_HOST_DEVICE void print() const
Definition: coordinate_transform.hpp:143
constexpr CK_TILE_HOST_DEVICE const auto & get_upper_lengths() const
Definition: coordinate_transform.hpp:87
static constexpr CK_TILE_HOST_DEVICE bool is_known_at_compile_time()
Definition: coordinate_transform.hpp:129
static constexpr CK_TILE_HOST_DEVICE auto get_type_enum()
Definition: coordinate_transform.hpp:82
static constexpr CK_TILE_HOST_DEVICE auto calculate_upper_dimension_safe_vector_length_strides(const LowVectorLengths &low_vector_lengths, const LowVectorStrides &low_vector_strides)
Definition: coordinate_transform.hpp:137
Definition: coordinate_transform.hpp:1110
static constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_always_mapped_to_valid_lower_index()
Definition: coordinate_transform.hpp:1138
static constexpr CK_TILE_HOST_DEVICE bool is_known_at_compile_time()
Definition: coordinate_transform.hpp:1150
constexpr CK_TILE_HOST_DEVICE replicate()=default
static constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_mapped_to_valid_lower_index(const UpIdx &)
Definition: coordinate_transform.hpp:1145
constexpr CK_TILE_HOST_DEVICE void calculate_lower_index(LowIdx &, const UpIdx &) const
Definition: coordinate_transform.hpp:1122
static CK_TILE_HOST_DEVICE void update_lower_index(LowIdxDiff &, const UpIdxDiff &, LowIdx &, const UpIdx &)
Definition: coordinate_transform.hpp:1130
CK_TILE_HOST_DEVICE void print() const
Definition: coordinate_transform.hpp:1155
constexpr CK_TILE_HOST_DEVICE auto get_upper_lengths() const
Definition: coordinate_transform.hpp:1119
constexpr CK_TILE_HOST_DEVICE replicate(const UpLengths &up_lengths)
Definition: coordinate_transform.hpp:1115
Definition: coordinate_transform.hpp:353
constexpr CK_TILE_HOST_DEVICE right_pad(const LowLength &low_length, const RightPadLength &right_pad_length)
Definition: coordinate_transform.hpp:365
constexpr CK_TILE_HOST_DEVICE const auto & get_upper_lengths() const
Definition: coordinate_transform.hpp:373
static constexpr CK_TILE_HOST_DEVICE bool is_known_at_compile_time()
Definition: coordinate_transform.hpp:415
static CK_TILE_HOST_DEVICE void update_lower_index(LowIdxDiff &idx_diff_low, const UpIdxDiff &idx_diff_up, LowIdx &idx_low, const UpIdx &)
Definition: coordinate_transform.hpp:386
static constexpr CK_TILE_HOST_DEVICE void calculate_lower_index(LowIdx &idx_low, const UpIdx &idx_up)
Definition: coordinate_transform.hpp:376
static constexpr CK_TILE_HOST_DEVICE auto calculate_upper_dimension_safe_vector_length_strides(const LowVectorLengths &low_vector_lengths, const LowVectorStrides &low_vector_strides)
Definition: coordinate_transform.hpp:425
RightPadLength right_pad_length_
Definition: coordinate_transform.hpp:361
CK_TILE_HOST_DEVICE void print() const
Definition: coordinate_transform.hpp:434
constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_mapped_to_valid_lower_index(const UpIdx &idx_up) const
Definition: coordinate_transform.hpp:410
static constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_always_mapped_to_valid_lower_index()
Definition: coordinate_transform.hpp:403
constexpr CK_TILE_HOST_DEVICE right_pad()=default
decltype(make_tuple(LowLength{}+RightPadLength{})) UpLengths
Definition: coordinate_transform.hpp:357
Definition: coordinate_transform.hpp:1172
constexpr CK_TILE_HOST_DEVICE slice(const LowLength &, const SliceBegin &slice_begin, const SliceEnd &slice_end)
Definition: coordinate_transform.hpp:1184
static CK_TILE_HOST_DEVICE void update_lower_index(LowIdxDiff &idx_diff_low, const UpIdxDiff &idx_diff_up, LowIdx &idx_low, const UpIdx &)
Definition: coordinate_transform.hpp:1206
constexpr CK_TILE_HOST_DEVICE slice()=default
constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_mapped_to_valid_lower_index(const UpIdx &) const
Definition: coordinate_transform.hpp:1230
decltype(make_tuple(SliceEnd{} - SliceBegin{})) UpLengths
Definition: coordinate_transform.hpp:1176
constexpr CK_TILE_HOST_DEVICE const auto & get_upper_lengths() const
Definition: coordinate_transform.hpp:1193
static constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_always_mapped_to_valid_lower_index()
Definition: coordinate_transform.hpp:1223
constexpr CK_TILE_HOST_DEVICE void calculate_lower_index(LowIdx &idx_low, const UpIdx &idx_up) const
Definition: coordinate_transform.hpp:1196
CK_TILE_HOST_DEVICE void print() const
Definition: coordinate_transform.hpp:1242
static constexpr CK_TILE_HOST_DEVICE bool is_known_at_compile_time()
Definition: coordinate_transform.hpp:1235
Definition: functional.hpp:43
Definition: tuple.hpp:192
Definition: coordinate_transform.hpp:858
static constexpr CK_TILE_HOST_DEVICE auto calculate_upper_dimension_safe_vector_length_strides(const LowVectorLengths &low_vector_lengths, const LowVectorStrides &low_vector_strides)
Definition: coordinate_transform.hpp:941
CK_TILE_HOST_DEVICE void update_lower_index(LowIdxDiff &idx_diff_low, const UpIdxDiff &idx_diff_up, LowIdx &idx_low, const UpIdx &) const
Definition: coordinate_transform.hpp:909
static constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_mapped_to_valid_lower_index(const UpIdx &)
Definition: coordinate_transform.hpp:927
constexpr CK_TILE_HOST_DEVICE unmerge(const UpLengths &up_lengths)
Definition: coordinate_transform.hpp:872
constexpr CK_TILE_HOST_DEVICE void calculate_lower_index(LowIdx &idx_low, const UpIdx &idx_up) const
Definition: coordinate_transform.hpp:886
static constexpr CK_TILE_HOST_DEVICE bool is_known_at_compile_time()
Definition: coordinate_transform.hpp:932
constexpr CK_TILE_HOST_DEVICE const auto & get_upper_lengths() const
Definition: coordinate_transform.hpp:883
UpLengthsScan up_lengths_scan_
Definition: coordinate_transform.hpp:868
static constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_always_mapped_to_valid_lower_index()
Definition: coordinate_transform.hpp:920
CK_TILE_HOST_DEVICE void print() const
Definition: coordinate_transform.hpp:962
constexpr CK_TILE_HOST_DEVICE unmerge()=default
static constexpr CK_TILE_HOST_DEVICE auto get_type_enum()
Definition: coordinate_transform.hpp:878
decltype(container_reverse_exclusive_scan(UpLengths{}, multiplies{}, number< 1 >{})) UpLengthsScan
Definition: coordinate_transform.hpp:865
Definition: coordinate_transform.hpp:1347
constexpr CK_TILE_HOST_DEVICE const auto & get_upper_lengths() const
Definition: coordinate_transform.hpp:1366
constexpr CK_TILE_HOST_DEVICE xor_t(const LowLengths &low_lengths)
Definition: coordinate_transform.hpp:1359
constexpr CK_TILE_HOST_DEVICE void calculate_lower_index(LowIdx &idx_low, const UpIdx &idx_up) const
Definition: coordinate_transform.hpp:1369
constexpr CK_TILE_HOST_DEVICE auto calculate_upper_dimension_safe_vector_length_strides(const LowVectorLengths &low_vector_lengths, const LowVectorStrides &low_vector_strides) const
Definition: coordinate_transform.hpp:1418
static constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_mapped_to_valid_lower_index(const UpIdx &)
Definition: coordinate_transform.hpp:1406
static constexpr CK_TILE_HOST_DEVICE bool is_known_at_compile_time()
Definition: coordinate_transform.hpp:1411
CK_TILE_HOST_DEVICE void print() const
Definition: coordinate_transform.hpp:1428
static constexpr CK_TILE_HOST_DEVICE auto get_type_enum()
Definition: coordinate_transform.hpp:1361
CK_TILE_HOST_DEVICE void update_lower_index(LowIdxDiff &idx_diff_low, const UpIdxDiff &, LowIdx &idx_low, const UpIdx &idx_up) const
Definition: coordinate_transform.hpp:1382
static constexpr CK_TILE_HOST_DEVICE bool is_valid_upper_index_always_mapped_to_valid_lower_index()
Definition: coordinate_transform.hpp:1399