detail Struct Reference

detail Struct Reference#

Composable Kernel: ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail Struct Reference
ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail Struct Reference

#include <tile_distribution_encoding.hpp>

Static Public Member Functions

static CK_TILE_HOST_DEVICE constexpr auto get_uniformed_h_dim_lengths ()
static CK_TILE_HOST_DEVICE constexpr auto get_uniformed_p_dim_lengths_over_h ()
static CK_TILE_HOST_DEVICE constexpr auto get_uniformed_rh_dim_lengths ()
static CK_TILE_HOST_DEVICE constexpr auto get_h_dim_lengths_prefix_sum ()
static CK_TILE_HOST_DEVICE constexpr auto get_rh_dim_lengths_prefix_sum ()
static CK_TILE_HOST_DEVICE constexpr auto get_uniformed_idx_p_to_h ()
static CK_TILE_HOST_DEVICE constexpr auto get_uniformed_idx_y_to_rh ()
static CK_TILE_HOST_DEVICE constexpr auto get_uniformed_idx_y_to_h ()
static CK_TILE_HOST_DEVICE constexpr auto get_y_to_h_masks ()
template<typename IdxSeq, typename PrefixSumSeq>
static CK_TILE_HOST_DEVICE constexpr auto get_sorted_info (IdxSeq, PrefixSumSeq)
static CK_TILE_HOST_DEVICE constexpr auto get_sorted_y_to_h_info ()

Static Public Attributes

static constexpr index_t ndim_rh_major_ = NDimX + 1
static constexpr index_t ndim_span_major_ = NDimX
static constexpr auto ndims_rhs_minor_
static constexpr index_t max_ndim_rh_minor_
static constexpr auto rhs_lengthss_
static constexpr auto ys_lengths_
static constexpr auto rhs_major_minor_to_ys_
static constexpr auto ndims_span_minor_
static constexpr index_t max_ndim_span_minor_
static constexpr auto rhs_major_minor_to_span_minor_
static constexpr auto ys_to_span_major_
static constexpr auto ys_to_span_minor_
static constexpr auto distributed_spans_lengthss_
static constexpr auto ndims_distributed_spans_minor_
static constexpr auto does_p_own_r_
static constexpr auto ps_over_rs_derivative_

Member Function Documentation

◆ get_h_dim_lengths_prefix_sum()

template<typename RsLengths_, typename HsLengthss_, typename Ps2RHssMajor_, typename Ps2RHssMinor_, typename Ys2RHsMajor_, typename Ys2RHsMinor_>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail::get_h_dim_lengths_prefix_sum ( )
inlinestaticconstexpr

◆ get_rh_dim_lengths_prefix_sum()

template<typename RsLengths_, typename HsLengthss_, typename Ps2RHssMajor_, typename Ps2RHssMinor_, typename Ys2RHsMajor_, typename Ys2RHsMinor_>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail::get_rh_dim_lengths_prefix_sum ( )
inlinestaticconstexpr

◆ get_sorted_info()

template<typename RsLengths_, typename HsLengthss_, typename Ps2RHssMajor_, typename Ps2RHssMinor_, typename Ys2RHsMajor_, typename Ys2RHsMinor_>
template<typename IdxSeq, typename PrefixSumSeq>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail::get_sorted_info ( IdxSeq ,
PrefixSumSeq  )
inlinestaticconstexpr

◆ get_sorted_y_to_h_info()

template<typename RsLengths_, typename HsLengthss_, typename Ps2RHssMajor_, typename Ps2RHssMinor_, typename Ys2RHsMajor_, typename Ys2RHsMinor_>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail::get_sorted_y_to_h_info ( )
inlinestaticconstexpr

◆ get_uniformed_h_dim_lengths()

template<typename RsLengths_, typename HsLengthss_, typename Ps2RHssMajor_, typename Ps2RHssMinor_, typename Ys2RHsMajor_, typename Ys2RHsMinor_>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail::get_uniformed_h_dim_lengths ( )
inlinestaticconstexpr

◆ get_uniformed_idx_p_to_h()

template<typename RsLengths_, typename HsLengthss_, typename Ps2RHssMajor_, typename Ps2RHssMinor_, typename Ys2RHsMajor_, typename Ys2RHsMinor_>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail::get_uniformed_idx_p_to_h ( )
inlinestaticconstexpr

◆ get_uniformed_idx_y_to_h()

template<typename RsLengths_, typename HsLengthss_, typename Ps2RHssMajor_, typename Ps2RHssMinor_, typename Ys2RHsMajor_, typename Ys2RHsMinor_>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail::get_uniformed_idx_y_to_h ( )
inlinestaticconstexpr

◆ get_uniformed_idx_y_to_rh()

template<typename RsLengths_, typename HsLengthss_, typename Ps2RHssMajor_, typename Ps2RHssMinor_, typename Ys2RHsMajor_, typename Ys2RHsMinor_>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail::get_uniformed_idx_y_to_rh ( )
inlinestaticconstexpr

◆ get_uniformed_p_dim_lengths_over_h()

template<typename RsLengths_, typename HsLengthss_, typename Ps2RHssMajor_, typename Ps2RHssMinor_, typename Ys2RHsMajor_, typename Ys2RHsMinor_>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail::get_uniformed_p_dim_lengths_over_h ( )
inlinestaticconstexpr

◆ get_uniformed_rh_dim_lengths()

template<typename RsLengths_, typename HsLengthss_, typename Ps2RHssMajor_, typename Ps2RHssMinor_, typename Ys2RHsMajor_, typename Ys2RHsMinor_>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail::get_uniformed_rh_dim_lengths ( )
inlinestaticconstexpr

◆ get_y_to_h_masks()

template<typename RsLengths_, typename HsLengthss_, typename Ps2RHssMajor_, typename Ps2RHssMinor_, typename Ys2RHsMajor_, typename Ys2RHsMinor_>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail::get_y_to_h_masks ( )
inlinestaticconstexpr

Member Data Documentation

◆ distributed_spans_lengthss_

template<typename RsLengths_, typename HsLengthss_, typename Ps2RHssMajor_, typename Ps2RHssMinor_, typename Ys2RHsMajor_, typename Ys2RHsMinor_>
auto ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail::distributed_spans_lengthss_
staticconstexpr
Initial value:
= [] {
distributed_spans_lengthss{{-1}};
static_for<0, NDimY, 1>{}([&](auto i) {
const index_t rh_major = ys_to_rhs_major_[i];
const index_t rh_minor = ys_to_rhs_minor_[i];
const index_t h_length = hs_lengthss_[number<rh_major - 1>{}][rh_minor];
const index_t span_major = rh_major - 1;
const index_t span_minor = rhs_major_minor_to_span_minor_[rh_major][rh_minor];
distributed_spans_lengthss(span_major)(span_minor) = h_length;
});
}()
constant< v > number
Definition tile/core/numeric/integral_constant.hpp:37
int32_t index_t
Definition integer.hpp:9
A fixed-size array container similar to std::array with additional utilities.
Definition tile/core/container/array.hpp:43
Definition tile/core/utility/functional.hpp:43
static constexpr auto rhs_major_minor_to_span_minor_
Definition tile_distribution_encoding.hpp:133
static constexpr index_t ndim_span_major_
Definition tile_distribution_encoding.hpp:61
Definition tile_distribution_encoding.hpp:26
static constexpr auto ys_to_rhs_major_
Definition tile_distribution_encoding.hpp:47
static constexpr auto ys_to_rhs_minor_
Definition tile_distribution_encoding.hpp:48
static constexpr auto hs_lengthss_
Definition tile_distribution_encoding.hpp:44

◆ does_p_own_r_

template<typename RsLengths_, typename HsLengthss_, typename Ps2RHssMajor_, typename Ps2RHssMinor_, typename Ys2RHsMajor_, typename Ys2RHsMinor_>
auto ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail::does_p_own_r_
staticconstexpr
Initial value:
= [] {
if constexpr(NDimR > 0)
{
array<array<bool, NDimR>, NDimP> does_p_own_r{{false}};
static_for<0, NDimP, 1>{}([&](auto idim_p) {
constexpr index_t ndim_low = ps_to_rhss_major_[idim_p].size();
static_for<0, ndim_low, 1>{}([&](auto idim_low) {
constexpr index_t rh_major = ps_to_rhss_major_[idim_p][idim_low];
constexpr index_t rh_minor = ps_to_rhss_minor_[idim_p][idim_low];
if constexpr(rh_major == 0)
{
does_p_own_r(idim_p)(rh_minor) = true;
}
});
});
return does_p_own_r;
}
else
{
}
}()
static constexpr index_t NDimR
Definition tile_distribution_encoding.hpp:40
static constexpr auto ps_to_rhss_minor_
Definition tile_distribution_encoding.hpp:46
static constexpr index_t NDimP
Definition tile_distribution_encoding.hpp:38
static constexpr auto ps_to_rhss_major_
Definition tile_distribution_encoding.hpp:45

◆ max_ndim_rh_minor_

template<typename RsLengths_, typename HsLengthss_, typename Ps2RHssMajor_, typename Ps2RHssMinor_, typename Ys2RHsMajor_, typename Ys2RHsMinor_>
index_t ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail::max_ndim_rh_minor_
staticconstexpr
Initial value:
=
CK_TILE_HOST_DEVICE constexpr auto container_reduce(const Container &x, Reduce reduce, Init init, number< IBegin >=number< 0 >{}, number< IEnd >=number< Container::size()>{}, number< IStep >=number< 1 >{})
Definition tile/core/container/container_helper.hpp:198
Definition tile/core/numeric/math.hpp:122
static constexpr auto ndims_rhs_minor_
Definition tile_distribution_encoding.hpp:64

◆ max_ndim_span_minor_

template<typename RsLengths_, typename HsLengthss_, typename Ps2RHssMajor_, typename Ps2RHssMinor_, typename Ys2RHsMajor_, typename Ys2RHsMinor_>
index_t ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail::max_ndim_span_minor_
staticconstexpr
Initial value:
=
static constexpr auto ndims_span_minor_
Definition tile_distribution_encoding.hpp:115

◆ ndim_rh_major_

template<typename RsLengths_, typename HsLengthss_, typename Ps2RHssMajor_, typename Ps2RHssMinor_, typename Ys2RHsMajor_, typename Ys2RHsMinor_>
index_t ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail::ndim_rh_major_ = NDimX + 1
staticconstexpr

◆ ndim_span_major_

template<typename RsLengths_, typename HsLengthss_, typename Ps2RHssMajor_, typename Ps2RHssMinor_, typename Ys2RHsMajor_, typename Ys2RHsMinor_>
index_t ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail::ndim_span_major_ = NDimX
staticconstexpr

◆ ndims_distributed_spans_minor_

template<typename RsLengths_, typename HsLengthss_, typename Ps2RHssMajor_, typename Ps2RHssMinor_, typename Ys2RHsMajor_, typename Ys2RHsMinor_>
auto ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail::ndims_distributed_spans_minor_
staticconstexpr
Initial value:
= [] {
array<index_t, ndim_span_major_> ndims_distributed_spans_minor{0};
static_for<0, NDimY, 1>{}([&](auto i) {
const index_t span_major = ys_to_rhs_major_[i] - 1;
ndims_distributed_spans_minor(span_major)++;
});
return ndims_distributed_spans_minor;
}()

◆ ndims_rhs_minor_

template<typename RsLengths_, typename HsLengthss_, typename Ps2RHssMajor_, typename Ps2RHssMinor_, typename Ys2RHsMajor_, typename Ys2RHsMinor_>
auto ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail::ndims_rhs_minor_
staticconstexpr
Initial value:
[](auto i) {
if constexpr(i.value == 0)
{
return rs_lengths_.size();
}
else
{
return hs_lengthss_[i - number<1>{}].size();
}
},
CK_TILE_HOST_DEVICE constexpr auto generate_array(F &&f, number< N >)
Definition tile/core/container/sequence.hpp:1115
static constexpr auto rs_lengths_
Definition tile_distribution_encoding.hpp:43

◆ ndims_span_minor_

template<typename RsLengths_, typename HsLengthss_, typename Ps2RHssMajor_, typename Ps2RHssMinor_, typename Ys2RHsMajor_, typename Ys2RHsMinor_>
auto ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail::ndims_span_minor_
staticconstexpr
Initial value:
= [] {
array<index_t, NDimX> ndims_span_minor{0};
for(index_t i = 0; i < NDimY; i++)
{
const index_t span_major = ys_to_rhs_major_[i] - 1;
ndims_span_minor(span_major)++;
}
return ndims_span_minor;
}()
static constexpr index_t NDimY
Definition tile_distribution_encoding.hpp:39

◆ ps_over_rs_derivative_

template<typename RsLengths_, typename HsLengthss_, typename Ps2RHssMajor_, typename Ps2RHssMinor_, typename Ys2RHsMajor_, typename Ys2RHsMinor_>
auto ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail::ps_over_rs_derivative_
staticconstexpr

◆ rhs_lengthss_

template<typename RsLengths_, typename HsLengthss_, typename Ps2RHssMajor_, typename Ps2RHssMinor_, typename Ys2RHsMajor_, typename Ys2RHsMinor_>
auto ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail::rhs_lengthss_
staticconstexpr
Initial value:
=
CK_TILE_HOST_DEVICE constexpr auto container_concat(const X &x, const Ys &... ys)
Definition tile/core/container/container_helper.hpp:363
CK_TILE_HOST_DEVICE constexpr auto to_array_of_array(tuple< Seqs... > t_of_s)
Definition tile/core/container/tuple.hpp:630
CK_TILE_HOST_DEVICE constexpr auto make_tuple(Xs &&... xs)
Definition tile/core/container/tuple.hpp:360

◆ rhs_major_minor_to_span_minor_

template<typename RsLengths_, typename HsLengthss_, typename Ps2RHssMajor_, typename Ps2RHssMinor_, typename Ys2RHsMajor_, typename Ys2RHsMinor_>
auto ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail::rhs_major_minor_to_span_minor_
staticconstexpr
Initial value:
= [] {
{-1}};
static_for<0, ndim_rh_major_, 1>{}([&](auto rh_major) {
constexpr index_t ndim_rh_minor = ndims_rhs_minor_[rh_major];
index_t cnt_ndim_span_minor = 0;
static_for<0, ndim_rh_minor, 1>{}([&](auto rh_minor) {
constexpr index_t idim_y = rhs_major_minor_to_ys_[rh_major][rh_minor];
if(idim_y >= 0)
{
rhs_major_minor_to_span_minor(rh_major)(rh_minor) = cnt_ndim_span_minor;
cnt_ndim_span_minor++;
}
});
});
}()
static constexpr index_t ndim_rh_major_
Definition tile_distribution_encoding.hpp:60
static constexpr auto rhs_major_minor_to_ys_
Definition tile_distribution_encoding.hpp:101

◆ rhs_major_minor_to_ys_

template<typename RsLengths_, typename HsLengthss_, typename Ps2RHssMajor_, typename Ps2RHssMinor_, typename Ys2RHsMajor_, typename Ys2RHsMinor_>
auto ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail::rhs_major_minor_to_ys_
staticconstexpr
Initial value:
= [] {
array<array<index_t, max_ndim_rh_minor_>, NDimX + 1> rhs_major_minor_to_ys_tmp{{-1}};
static_for<0, NDimY, 1>{}([&](auto i) {
constexpr index_t rh_major = ys_to_rhs_major_[i];
constexpr index_t rh_minor = ys_to_rhs_minor_[i];
rhs_major_minor_to_ys_tmp(rh_major)(rh_minor) = i;
});
return rhs_major_minor_to_ys_tmp;
}()
static constexpr index_t NDimX
Definition tile_distribution_encoding.hpp:37

◆ ys_lengths_

template<typename RsLengths_, typename HsLengthss_, typename Ps2RHssMajor_, typename Ps2RHssMinor_, typename Ys2RHsMajor_, typename Ys2RHsMinor_>
auto ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail::ys_lengths_
staticconstexpr
Initial value:
= [] {
array<index_t, NDimY> ys_lengths_tmp{-1};
for(index_t i = 0; i < NDimY; i++)
{
index_t rh_major = ys_to_rhs_major_[i];
index_t rh_minor = ys_to_rhs_minor_[i];
ys_lengths_tmp(i) = rhs_lengthss_[rh_major][rh_minor];
}
return ys_lengths_tmp;
}()
static constexpr auto rhs_lengthss_
Definition tile_distribution_encoding.hpp:82

◆ ys_to_span_major_

template<typename RsLengths_, typename HsLengthss_, typename Ps2RHssMajor_, typename Ps2RHssMinor_, typename Ys2RHsMajor_, typename Ys2RHsMinor_>
auto ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail::ys_to_span_major_
staticconstexpr
Initial value:
=
generate_array([](auto i) { return ys_to_rhs_major_[i] - 1; }, number<NDimY>{})

◆ ys_to_span_minor_

template<typename RsLengths_, typename HsLengthss_, typename Ps2RHssMajor_, typename Ps2RHssMinor_, typename Ys2RHsMajor_, typename Ys2RHsMinor_>
auto ck_tile::tile_distribution_encoding< RsLengths_, HsLengthss_, Ps2RHssMajor_, Ps2RHssMinor_, Ys2RHsMajor_, Ys2RHsMinor_ >::detail::ys_to_span_minor_
staticconstexpr

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