TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection > Struct Template Reference

TransposeTileDistributionTraits&lt; TileDistributionEncoding_, DataType_, Policy, ReverseDirection &gt; Struct Template Reference#

Composable Kernel: ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection > Struct Template Reference
ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection > Struct Template Reference

#include <load_tile_transpose.hpp>

Public Types

using InDstrEncode = remove_cvref_t<TileDistributionEncoding_>
using QuadInputEncoding
using QuadOutputEncoding
using TransposedDstrEncode

Static Public Attributes

static constexpr auto input_hs_lengthss = InDstrEncode::hs_lengthss_
static constexpr index_t LaneGroupSize
static constexpr auto quad_input_hs_lengthss = QuadInputEncoding::hs_lengthss_
static constexpr auto quad_output_hs_lengthss = QuadOutputEncoding::hs_lengthss_
static constexpr auto input_ps_to_rhss_major = InDstrEncode::ps_to_rhss_major_
static constexpr auto input_ps_to_rhss_minor = InDstrEncode::ps_to_rhss_minor_
static constexpr auto input_ys_to_rhs_major = InDstrEncode::ys_to_rhs_major_
static constexpr auto input_ys_to_rhs_minor = InDstrEncode::ys_to_rhs_minor_
static constexpr auto I0 = number<0>{}
static constexpr auto quad_input_ps_to_rhss_major0 = QuadInputEncoding::ps_to_rhss_major_[I0]
static constexpr auto quad_input_ps_to_rhss_minor0 = QuadInputEncoding::ps_to_rhss_minor_[I0]
static constexpr auto quad_output_ps_to_rhss_major0 = QuadOutputEncoding::ps_to_rhss_major_[I0]
static constexpr auto quad_output_ps_to_rhss_minor0 = QuadOutputEncoding::ps_to_rhss_minor_[I0]
static constexpr auto quad_output_ys_to_rhs_major = QuadOutputEncoding::ys_to_rhs_major_
static constexpr auto quad_output_ys_to_rhs_minor = QuadOutputEncoding::ys_to_rhs_minor_
static constexpr index_t dim0 = Policy::transpose_dims[0]
static constexpr index_t dim1 = Policy::transpose_dims[1]
static constexpr auto swap_one_and_two
static constexpr auto outer_hs_lengthss
static constexpr auto reversed_outer_hs_lengthss = tuple_reverse(outer_hs_lengthss)
static constexpr auto dst_out_hs_lengthss
static constexpr auto dst_ps_to_rhss_major
static constexpr auto quad_idx_offset
static constexpr auto quad_output_ps_minor_offset
static constexpr auto quad_output_ys_minor_offset
static constexpr auto dst_ps_to_rhss_minor
static constexpr auto outer_input_ys_to_rhs_major = input_ys_to_rhs_major.pop_back()
static constexpr auto dst_ys_to_rhs_major
static constexpr auto dst_ys_to_rhs_minor

Member Typedef Documentation

◆ InDstrEncode

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
using ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::InDstrEncode = remove_cvref_t<TileDistributionEncoding_>

◆ QuadInputEncoding

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
using ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::QuadInputEncoding
Initial value:
std::conditional_t<
ReverseDirection,
typename Policy::template QuadOutputEncoding<LaneGroupSize>,
typename Policy::template QuadInputEncoding<LaneGroupSize>>
std::conditional_t< ReverseDirection, typename Policy::template QuadInputEncoding< LaneGroupSize >, typename Policy::template QuadOutputEncoding< LaneGroupSize > > QuadOutputEncoding
Definition load_tile_transpose.hpp:219
std::conditional_t< ReverseDirection, typename Policy::template QuadOutputEncoding< LaneGroupSize >, typename Policy::template QuadInputEncoding< LaneGroupSize > > QuadInputEncoding
Definition load_tile_transpose.hpp:215

◆ QuadOutputEncoding

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
using ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::QuadOutputEncoding
Initial value:
std::conditional_t<
ReverseDirection,
typename Policy::template QuadInputEncoding<LaneGroupSize>,
typename Policy::template QuadOutputEncoding<LaneGroupSize>>

◆ TransposedDstrEncode

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
using ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::TransposedDstrEncode
Initial value:
tile_distribution_encoding<typename InDstrEncode::RsLengths,
remove_cv_t< std::remove_reference_t< T > > remove_cvref_t
Definition type_traits.hpp:21
static constexpr auto dst_ps_to_rhss_major
Definition load_tile_transpose.hpp:268
static constexpr auto dst_ys_to_rhs_minor
Definition load_tile_transpose.hpp:323
static constexpr auto dst_out_hs_lengthss
Definition load_tile_transpose.hpp:258
static constexpr auto dst_ps_to_rhss_minor
Definition load_tile_transpose.hpp:297
static constexpr auto dst_ys_to_rhs_major
Definition load_tile_transpose.hpp:320
Definition tile_distribution_encoding.hpp:26

Member Data Documentation

◆ dim0

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
index_t ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::dim0 = Policy::transpose_dims[0]
staticconstexpr

◆ dim1

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
index_t ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::dim1 = Policy::transpose_dims[1]
staticconstexpr

◆ dst_out_hs_lengthss

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
auto ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::dst_out_hs_lengthss
staticconstexpr
Initial value:
[](auto i) {
auto outer_i = reversed_outer_hs_lengthss[i];
return outer_i.push_back(quad_output_hs_lengthss[i]);
},
constant< v > number
Definition tile/core/numeric/integral_constant.hpp:37
CK_TILE_HOST_DEVICE constexpr auto generate_tuple(F &&f, number< N >)
Definition tile/core/container/tuple.hpp:429
static constexpr auto reversed_outer_hs_lengthss
Definition load_tile_transpose.hpp:257
static constexpr auto quad_output_hs_lengthss
Definition load_tile_transpose.hpp:225

◆ dst_ps_to_rhss_major

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
auto ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::dst_ps_to_rhss_major
staticconstexpr
Initial value:
[](auto i) {
if constexpr(i == input_ps_to_rhss_major.size() - 1)
{
constexpr auto current_size = input_ps_to_rhss_major[i].size();
constexpr auto reduce_size = quad_input_ps_to_rhss_major0.size();
constexpr auto quad_out = quad_output_ps_to_rhss_major0;
constexpr auto reduced_ps_to_rhss_major = input_ps_to_rhss_major[i].extract(
return reduced_ps_to_rhss_major.transform(swap_one_and_two).push_back(quad_out);
}
else
{
}
},
Definition load_tile_transpose.hpp:207
static constexpr auto quad_input_ps_to_rhss_major0
Definition load_tile_transpose.hpp:233
static constexpr auto quad_output_ps_to_rhss_major0
Definition load_tile_transpose.hpp:235
static constexpr auto swap_one_and_two
Definition load_tile_transpose.hpp:243
static constexpr auto input_ps_to_rhss_major
Definition load_tile_transpose.hpp:227
typename std::conditional< kHasContent, type0, type1 >::type type
Definition tile/core/container/sequence.hpp:302

◆ dst_ps_to_rhss_minor

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
auto ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::dst_ps_to_rhss_minor
staticconstexpr
Initial value:
[](auto i) {
constexpr auto input_i = input_ps_to_rhss_minor[i];
if constexpr(i == input_ps_to_rhss_minor.size() - 1)
{
constexpr auto outer_len = input_i.size() - quad_input_ps_to_rhss_minor0.size();
constexpr auto outer_ps =
typename sequence_split<decltype(input_i), outer_len>::left_type{};
return outer_ps.push_back(quad_output_ps_minor_offset +
}
else
{
return input_i;
}
},
static constexpr auto quad_output_ps_to_rhss_minor0
Definition load_tile_transpose.hpp:236
static constexpr auto quad_input_ps_to_rhss_minor0
Definition load_tile_transpose.hpp:234
static constexpr auto quad_output_ps_minor_offset
Definition load_tile_transpose.hpp:292
static constexpr auto input_ps_to_rhss_minor
Definition load_tile_transpose.hpp:228
Definition tile/core/container/sequence.hpp:352

◆ dst_ys_to_rhs_major

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
auto ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::dst_ys_to_rhs_major
staticconstexpr
Initial value:
=
static constexpr auto outer_input_ys_to_rhs_major
Definition load_tile_transpose.hpp:317

◆ dst_ys_to_rhs_minor

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
auto ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::dst_ys_to_rhs_minor
staticconstexpr
Initial value:
= input_ys_to_rhs_minor.pop_back().push_back(
static constexpr auto quad_output_ys_to_rhs_minor
Definition load_tile_transpose.hpp:238
static constexpr auto I0
Definition load_tile_transpose.hpp:232
static constexpr auto input_ys_to_rhs_minor
Definition load_tile_transpose.hpp:230
static constexpr auto quad_output_ys_minor_offset
Definition load_tile_transpose.hpp:294

◆ I0

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
auto ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::I0 = number<0>{}
staticconstexpr

◆ input_hs_lengthss

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
auto ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::input_hs_lengthss = InDstrEncode::hs_lengthss_
staticconstexpr

◆ input_ps_to_rhss_major

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
auto ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::input_ps_to_rhss_major = InDstrEncode::ps_to_rhss_major_
staticconstexpr

◆ input_ps_to_rhss_minor

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
auto ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::input_ps_to_rhss_minor = InDstrEncode::ps_to_rhss_minor_
staticconstexpr

◆ input_ys_to_rhs_major

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
auto ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::input_ys_to_rhs_major = InDstrEncode::ys_to_rhs_major_
staticconstexpr

◆ input_ys_to_rhs_minor

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
auto ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::input_ys_to_rhs_minor = InDstrEncode::ys_to_rhs_minor_
staticconstexpr

◆ LaneGroupSize

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
index_t ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::LaneGroupSize
staticconstexpr
Initial value:
=
Policy::template ValidationTraits<InDstrEncode, ReverseDirection>::LaneGroupSize

◆ outer_hs_lengthss

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
auto ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::outer_hs_lengthss
staticconstexpr
Initial value:
[](auto i) {
constexpr auto input_i = input_hs_lengthss[i];
constexpr auto outer_len = input_i.size() - quad_input_hs_lengthss[i].size();
return typename sequence_split<decltype(input_i), outer_len>::left_type{};
},
static constexpr auto quad_input_hs_lengthss
Definition load_tile_transpose.hpp:224
static constexpr auto input_hs_lengthss
Definition load_tile_transpose.hpp:209

◆ outer_input_ys_to_rhs_major

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
auto ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::outer_input_ys_to_rhs_major = input_ys_to_rhs_major.pop_back()
staticconstexpr

◆ quad_idx_offset

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
auto ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::quad_idx_offset
staticconstexpr
Initial value:
=
transform_tuples([](auto x) { return number<x.size()>{}; }, reversed_outer_hs_lengthss)
CK_TILE_HOST_DEVICE constexpr auto transform_tuples(F f, const X &x)
Definition tile/core/container/tuple.hpp:505

◆ quad_input_hs_lengthss

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
auto ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::quad_input_hs_lengthss = QuadInputEncoding::hs_lengthss_
staticconstexpr

◆ quad_input_ps_to_rhss_major0

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
auto ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::quad_input_ps_to_rhss_major0 = QuadInputEncoding::ps_to_rhss_major_[I0]
staticconstexpr

◆ quad_input_ps_to_rhss_minor0

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
auto ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::quad_input_ps_to_rhss_minor0 = QuadInputEncoding::ps_to_rhss_minor_[I0]
staticconstexpr

◆ quad_output_hs_lengthss

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
auto ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::quad_output_hs_lengthss = QuadOutputEncoding::hs_lengthss_
staticconstexpr

◆ quad_output_ps_minor_offset

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
auto ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::quad_output_ps_minor_offset
staticconstexpr
Initial value:
[](auto x) { return quad_idx_offset[number<x - 1>{}]; }, quad_output_ps_to_rhss_major0))
CK_TILE_HOST_DEVICE constexpr auto generate_tuple_for(F &&f, sequence< ids... >)
Definition tile/core/container/tuple.hpp:423
CK_TILE_HOST_DEVICE constexpr auto to_sequence(tuple< number< Is >... >)
Definition tile/core/container/sequence.hpp:1055
static constexpr auto quad_idx_offset
Definition load_tile_transpose.hpp:288

◆ quad_output_ps_to_rhss_major0

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
auto ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::quad_output_ps_to_rhss_major0 = QuadOutputEncoding::ps_to_rhss_major_[I0]
staticconstexpr

◆ quad_output_ps_to_rhss_minor0

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
auto ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::quad_output_ps_to_rhss_minor0 = QuadOutputEncoding::ps_to_rhss_minor_[I0]
staticconstexpr

◆ quad_output_ys_minor_offset

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
auto ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::quad_output_ys_minor_offset
staticconstexpr
Initial value:
[](auto x) { return quad_idx_offset[number<x - 1>{}]; }, quad_output_ys_to_rhs_major))
static constexpr auto quad_output_ys_to_rhs_major
Definition load_tile_transpose.hpp:237

◆ quad_output_ys_to_rhs_major

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
auto ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::quad_output_ys_to_rhs_major = QuadOutputEncoding::ys_to_rhs_major_
staticconstexpr

◆ quad_output_ys_to_rhs_minor

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
auto ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::quad_output_ys_to_rhs_minor = QuadOutputEncoding::ys_to_rhs_minor_
staticconstexpr

◆ reversed_outer_hs_lengthss

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
auto ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::reversed_outer_hs_lengthss = tuple_reverse(outer_hs_lengthss)
staticconstexpr

◆ swap_one_and_two

template<typename TileDistributionEncoding_, typename DataType_, typename Policy = DefaultTranspose<DataType_>, bool ReverseDirection = false>
auto ck_tile::TransposeTileDistributionTraits< TileDistributionEncoding_, DataType_, Policy, ReverseDirection >::swap_one_and_two
staticconstexpr
Initial value:
= [](const index_t idx) {
return (idx == 1) ? 2 : (idx == 2) ? 1 : idx;
}
int32_t index_t
Definition integer.hpp:9

The documentation for this struct was generated from the following file: