space_filling_curve< TensorLengths, DimAccessOrder, ScalarsPerAccess, SnakeCurved > Struct Template Reference

space_filling_curve&lt; TensorLengths, DimAccessOrder, ScalarsPerAccess, SnakeCurved &gt; Struct Template Reference#

Composable Kernel: ck_tile::space_filling_curve< TensorLengths, DimAccessOrder, ScalarsPerAccess, SnakeCurved > Struct Template Reference
Public Types | Static Public Member Functions | Static Public Attributes | List of all members
ck_tile::space_filling_curve< TensorLengths, DimAccessOrder, ScalarsPerAccess, SnakeCurved > Struct Template Reference

#include <space_filling_curve.hpp>

Public Types

using Index = multi_index< nDim >
 

Static Public Member Functions

static constexpr CK_TILE_HOST_DEVICE index_t get_num_of_access ()
 
template<index_t AccessIdx1dHead, index_t AccessIdx1dTail>
static constexpr CK_TILE_HOST_DEVICE auto get_step_between (number< AccessIdx1dHead >, number< AccessIdx1dTail >)
 
template<index_t AccessIdx1d>
static constexpr CK_TILE_HOST_DEVICE auto get_forward_step (number< AccessIdx1d >)
 
template<index_t AccessIdx1d>
static constexpr CK_TILE_HOST_DEVICE auto get_backward_step (number< AccessIdx1d >)
 
template<index_t AccessIdx1d>
static constexpr CK_TILE_HOST_DEVICE Index _get_index (number< AccessIdx1d >)
 
template<index_t AccessIdx1d>
static constexpr CK_TILE_HOST_DEVICE auto get_index (number< AccessIdx1d >)
 

Static Public Attributes

static constexpr index_t TensorSize
 
static constexpr index_t nDim = TensorLengths::size()
 
static constexpr index_t ScalarPerVector
 
static constexpr auto access_lengths = TensorLengths{} / ScalarsPerAccess{}
 
static constexpr auto dim_access_order = DimAccessOrder{}
 
static constexpr auto ordered_access_lengths
 
static constexpr auto to_index_adaptor
 
static constexpr auto I0 = number<0>{}
 
static constexpr auto I1 = number<1>{}
 

Member Typedef Documentation

◆ Index

template<typename TensorLengths , typename DimAccessOrder , typename ScalarsPerAccess , bool SnakeCurved = true>
using ck_tile::space_filling_curve< TensorLengths, DimAccessOrder, ScalarsPerAccess, SnakeCurved >::Index = multi_index<nDim>

Member Function Documentation

◆ _get_index()

template<typename TensorLengths , typename DimAccessOrder , typename ScalarsPerAccess , bool SnakeCurved = true>
template<index_t AccessIdx1d>
static constexpr CK_TILE_HOST_DEVICE Index ck_tile::space_filling_curve< TensorLengths, DimAccessOrder, ScalarsPerAccess, SnakeCurved >::_get_index ( number< AccessIdx1d >  )
inlinestaticconstexpr

◆ get_backward_step()

template<typename TensorLengths , typename DimAccessOrder , typename ScalarsPerAccess , bool SnakeCurved = true>
template<index_t AccessIdx1d>
static constexpr CK_TILE_HOST_DEVICE auto ck_tile::space_filling_curve< TensorLengths, DimAccessOrder, ScalarsPerAccess, SnakeCurved >::get_backward_step ( number< AccessIdx1d >  )
inlinestaticconstexpr

◆ get_forward_step()

template<typename TensorLengths , typename DimAccessOrder , typename ScalarsPerAccess , bool SnakeCurved = true>
template<index_t AccessIdx1d>
static constexpr CK_TILE_HOST_DEVICE auto ck_tile::space_filling_curve< TensorLengths, DimAccessOrder, ScalarsPerAccess, SnakeCurved >::get_forward_step ( number< AccessIdx1d >  )
inlinestaticconstexpr

◆ get_index()

template<typename TensorLengths , typename DimAccessOrder , typename ScalarsPerAccess , bool SnakeCurved = true>
template<index_t AccessIdx1d>
static constexpr CK_TILE_HOST_DEVICE auto ck_tile::space_filling_curve< TensorLengths, DimAccessOrder, ScalarsPerAccess, SnakeCurved >::get_index ( number< AccessIdx1d >  )
inlinestaticconstexpr

◆ get_num_of_access()

template<typename TensorLengths , typename DimAccessOrder , typename ScalarsPerAccess , bool SnakeCurved = true>
static constexpr CK_TILE_HOST_DEVICE index_t ck_tile::space_filling_curve< TensorLengths, DimAccessOrder, ScalarsPerAccess, SnakeCurved >::get_num_of_access ( )
inlinestaticconstexpr

◆ get_step_between()

template<typename TensorLengths , typename DimAccessOrder , typename ScalarsPerAccess , bool SnakeCurved = true>
template<index_t AccessIdx1dHead, index_t AccessIdx1dTail>
static constexpr CK_TILE_HOST_DEVICE auto ck_tile::space_filling_curve< TensorLengths, DimAccessOrder, ScalarsPerAccess, SnakeCurved >::get_step_between ( number< AccessIdx1dHead >  ,
number< AccessIdx1dTail >   
)
inlinestaticconstexpr

Member Data Documentation

◆ access_lengths

template<typename TensorLengths , typename DimAccessOrder , typename ScalarsPerAccess , bool SnakeCurved = true>
constexpr auto ck_tile::space_filling_curve< TensorLengths, DimAccessOrder, ScalarsPerAccess, SnakeCurved >::access_lengths = TensorLengths{} / ScalarsPerAccess{}
staticconstexpr

◆ dim_access_order

template<typename TensorLengths , typename DimAccessOrder , typename ScalarsPerAccess , bool SnakeCurved = true>
constexpr auto ck_tile::space_filling_curve< TensorLengths, DimAccessOrder, ScalarsPerAccess, SnakeCurved >::dim_access_order = DimAccessOrder{}
staticconstexpr

◆ I0

template<typename TensorLengths , typename DimAccessOrder , typename ScalarsPerAccess , bool SnakeCurved = true>
constexpr auto ck_tile::space_filling_curve< TensorLengths, DimAccessOrder, ScalarsPerAccess, SnakeCurved >::I0 = number<0>{}
staticconstexpr

◆ I1

template<typename TensorLengths , typename DimAccessOrder , typename ScalarsPerAccess , bool SnakeCurved = true>
constexpr auto ck_tile::space_filling_curve< TensorLengths, DimAccessOrder, ScalarsPerAccess, SnakeCurved >::I1 = number<1>{}
staticconstexpr

◆ nDim

template<typename TensorLengths , typename DimAccessOrder , typename ScalarsPerAccess , bool SnakeCurved = true>
constexpr index_t ck_tile::space_filling_curve< TensorLengths, DimAccessOrder, ScalarsPerAccess, SnakeCurved >::nDim = TensorLengths::size()
staticconstexpr

◆ ordered_access_lengths

template<typename TensorLengths , typename DimAccessOrder , typename ScalarsPerAccess , bool SnakeCurved = true>
constexpr auto ck_tile::space_filling_curve< TensorLengths, DimAccessOrder, ScalarsPerAccess, SnakeCurved >::ordered_access_lengths
staticconstexpr
Initial value:
=
container_reorder_given_new2old(access_lengths, dim_access_order)
ck_tile::container_reorder_given_new2old
constexpr CK_TILE_HOST_DEVICE auto container_reorder_given_new2old(const array< TData, NSize > &old_array, sequence< IRs... >)
Definition: container_helper.hpp:39
ck_tile::space_filling_curve::dim_access_order
static constexpr auto dim_access_order
Definition: space_filling_curve.hpp:34
ck_tile::space_filling_curve::access_lengths
static constexpr auto access_lengths
Definition: space_filling_curve.hpp:33

◆ ScalarPerVector

template<typename TensorLengths , typename DimAccessOrder , typename ScalarsPerAccess , bool SnakeCurved = true>
constexpr index_t ck_tile::space_filling_curve< TensorLengths, DimAccessOrder, ScalarsPerAccess, SnakeCurved >::ScalarPerVector
staticconstexpr
Initial value:
=
reduce_on_sequence(ScalarsPerAccess{}, multiplies{}, number<1>{})
ck_tile::multiplies
__host__ __device__ multiplies() -> multiplies< void, void >
FIXME: create macro to replace 'host device' and nothing more.
ck_tile::reduce_on_sequence
constexpr CK_TILE_HOST_DEVICE index_t reduce_on_sequence(Seq, Reduce f, number< Init >)
Definition: sequence.hpp:973

◆ TensorSize

template<typename TensorLengths , typename DimAccessOrder , typename ScalarsPerAccess , bool SnakeCurved = true>
constexpr index_t ck_tile::space_filling_curve< TensorLengths, DimAccessOrder, ScalarsPerAccess, SnakeCurved >::TensorSize
staticconstexpr
Initial value:
=
reduce_on_sequence(TensorLengths{}, multiplies{}, number<1>{})

◆ to_index_adaptor

template<typename TensorLengths , typename DimAccessOrder , typename ScalarsPerAccess , bool SnakeCurved = true>
constexpr auto ck_tile::space_filling_curve< TensorLengths, DimAccessOrder, ScalarsPerAccess, SnakeCurved >::to_index_adaptor
staticconstexpr
Initial value:
= make_single_stage_tensor_adaptor(
make_tuple(make_merge_transform(ordered_access_lengths)),
make_tuple(typename arithmetic_sequence_gen<0, nDim, 1>::type{}),
make_tuple(sequence<0>{}))
ck_tile::make_merge_transform
constexpr CK_TILE_HOST_DEVICE auto make_merge_transform(const LowLengths &low_lengths)
Definition: coordinate_transform.hpp:1672
ck_tile::make_single_stage_tensor_adaptor
constexpr CK_TILE_HOST_DEVICE auto make_single_stage_tensor_adaptor(const Transforms &transforms, LowerDimensionOldTopIdss, UpperDimensionNewTopIdss)
Definition: tensor_adaptor.hpp:348
ck_tile::make_tuple
constexpr CK_TILE_HOST_DEVICE auto make_tuple(Xs &&... xs)
Definition: tuple.hpp:337
ck_tile::arithmetic_sequence_gen::type
typename std::conditional< kHasContent, type0, type1 >::type type
Definition: sequence.hpp:293
ck_tile::space_filling_curve::ordered_access_lengths
static constexpr auto ordered_access_lengths
Definition: space_filling_curve.hpp:35
The documentation for this struct was generated from the following file:
  • /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-6.4.3/include/ck_tile/core/algorithm/space_filling_curve.hpp