tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ > Struct Template Reference

• ck_tile
• tile_window_linear

Classes | Public Types | Public Member Functions | Static Public Member Functions | Public Attributes | Static Public Attributes | List of all members

ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ > Struct Template Reference

#include <tile_window_linear.hpp>

Inheritance diagram for ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >:

Classes
struct	traits

Public Types
using	Base = tile_window_with_tile_dstr_base< tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >, BottomTensorView_, WindowLengths_, StaticTileDistribution_ >
using	LinearBottomDims = remove_cvref_t< LinearBottomDims_ >
using	AccessMap_NonLinear = typename traits::AccessMap_NonLinear
using	AccessHistogram_NonLinear = typename traits::AccessHistogram_NonLinear
using	AccessPrefixSum_NonLinear = typename traits::AccessPrefixSum_NonLinear
Public Types inherited from ck_tile::tile_window_with_tile_dstr_base< tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >, BottomTensorView_, WindowLengths_, StaticTileDistribution_ >
using	TileDstr = remove_cvref_t< StaticTileDistribution_ >
using	TileWindowBase = tile_window_base< tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >, BottomTensorView_, WindowLengths_ >
using	WindowAdaptor = typename TileDstr::PsYs2XsAdaptor
using	AdaptorTopIndex = array< index_t, NDimWindowAdaptorTop >
using	WindowAdaptorCoord = decltype(make_tensor_adaptor_coordinate(WindowAdaptor{}, AdaptorTopIndex{}))
using	BottomTensorCoord = decltype(make_tensor_coordinate(typename TileWindowBase::BottomTensorDesc{}, typename TileWindowBase::BottomTensorIndex{}))
Public Types inherited from ck_tile::tile_window_base< tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >, BottomTensorView_, WindowLengths_ >
using	BottomTensorView = remove_reference_t< BottomTensorView_ >
using	WindowLengths = remove_cvref_t< WindowLengths_ >
using	BottomTensorDesc = typename BottomTensorView::TensorDesc
using	DataType = remove_cvref_t< typename BottomTensorView::DataType >
using	BottomTensorIndex = array< index_t, NDimBottomTensor >

Public Member Functions
constexpr CK_TILE_DEVICE	tile_window_linear ()=default
constexpr CK_TILE_DEVICE	tile_window_linear (const typename Base::BottomTensorView &bottom_tensor_view, const typename Base::WindowLengths &window_lengths, const typename Base::BottomTensorIndex &window_origin, const typename Base::TileDstr &tile_distribution)
template<index_t i_access = -1, bool oob_conditional_check = true>
CK_TILE_DEVICE auto	load (number< i_access >={}, bool_constant< oob_conditional_check >={}) const
template<typename DstTile , index_t i_access = -1, bool oob_conditional_check = true>
CK_TILE_DEVICE auto	load (DstTile &dst_tensor, number< i_access >={}, bool_constant< oob_conditional_check >={}) const
template<typename DstTile , index_t i_access = -1, bool oob_conditional_check = true, bool pre_nop = false>
CK_TILE_DEVICE void	load_raw (DstTile &dst_tensor, number< i_access >={}, bool_constant< oob_conditional_check >={}, bool_constant< pre_nop >={}) const
template<typename LdsTileWindow_ , index_t i_access = -1, bool oob_conditional_check = true, bool pre_nop = false>
CK_TILE_DEVICE auto	async_load_raw (LdsTileWindow_ &&lds_tile, number< i_access >={}, bool_constant< oob_conditional_check >={}, bool_constant< pre_nop >={}) const
template<typename LdsTileWindow_ , index_t i_access = -1, bool oob_conditional_check = true>
CK_TILE_DEVICE auto	async_load (LdsTileWindow_ &&lds_tile, number< i_access >={}, bool_constant< oob_conditional_check >={}) const
template<typename Policy , index_t i_access_unsupport_ = -1, bool oob_conditional_check = true>
CK_TILE_DEVICE auto	load_transpose () const
template<typename Policy , typename DistributedTensor , index_t i_access = -1, bool oob_conditional_check = true>
CK_TILE_DEVICE auto	load_transpose_linear (DistributedTensor &dst_tensor, number< i_access >={}, bool_constant< oob_conditional_check >={}) const
template<index_t i_access = -1, bool oob_conditional_check = true>
CK_TILE_DEVICE void	store (const static_distributed_tensor< typename Base::DataType, typename Base::TileDstr > &dstr_tensor, number< i_access >={}, bool_constant< oob_conditional_check >={}) const
template<index_t i_access = -1>
CK_TILE_DEVICE void	store_raw (const static_distributed_tensor< typename Base::DataType, typename Base::TileDstr > &dstr_tensor, number< i_access >={}) const
template<index_t i_access = -1, bool oob_conditional_check = true>
CK_TILE_DEVICE void	update (const static_distributed_tensor< typename Base::DataType, typename Base::TileDstr > &dstr_tensor, number< i_access >={}, bool_constant< oob_conditional_check >={}) const
template<index_t i_access = -1, bool oob_conditional_check = true, bool pre_nop = false>
CK_TILE_DEVICE void	update_raw (const static_distributed_tensor< typename Base::DataType, typename Base::TileDstr > &dstr_tensor, number< i_access >={}, bool_constant< oob_conditional_check >={}, bool_constant< pre_nop >={}) const
CK_TILE_DEVICE void	move_extended (const typename Base::BottomTensorIndex &step)
CK_TILE_DEVICE void	set_window_origin_extended (const typename Base::BottomTensorIndex &)
Public Member Functions inherited from ck_tile::tile_window_with_tile_dstr_base< tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >, BottomTensorView_, WindowLengths_, StaticTileDistribution_ >
constexpr CK_TILE_DEVICE auto	get_tile_distribution () const
CK_TILE_HOST_DEVICE void	init_raw ()
CK_TILE_DEVICE void	move_window_adaptor_and_bottom_tensor_thread_coordinate (WindowAdaptorCoord &window_adaptor_thread_coord, BottomTensorCoord &bottom_tensor_thread_coord, const ATopIndex &idx_diff_adaptor_top) const
constexpr CK_TILE_DEVICE auto	get_num_of_access () const
Public Member Functions inherited from ck_tile::tile_window_base< tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >, BottomTensorView_, WindowLengths_ >
constexpr CK_TILE_DEVICE auto	get_window_origin () const
constexpr CK_TILE_DEVICE auto	get_window_lengths () const
constexpr CK_TILE_DEVICE auto	get_bottom_tensor_view () const
CK_TILE_DEVICE void	set_window_origin (const BottomTensorIndex &new_window_origin)
CK_TILE_DEVICE void	set_window_origin_extended (const BottomTensorIndex &)
constexpr CK_TILE_DEVICE void	set_bottom_tensor_view_data_ptr (typename BottomTensorView::DataType *data)
CK_TILE_DEVICE void	move (const BottomTensorIndex &step)
CK_TILE_DEVICE void	move_extended (const BottomTensorIndex &)

Static Public Member Functions
template<index_t i_access>
static constexpr CK_TILE_DEVICE auto	get_bottom_linear_coordinate (number< i_access >)
template<index_t i_access>
static constexpr CK_TILE_DEVICE index_t	get_bottom_linear_offset (number< i_access >)
Static Public Member Functions inherited from ck_tile::tile_window_with_tile_dstr_base< tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >, BottomTensorView_, WindowLengths_, StaticTileDistribution_ >
static constexpr CK_TILE_DEVICE bool	has_static_tile_distribution ()
static constexpr CK_TILE_DEVICE auto	get_window_adaptor_ys_safe_vector_length_strides ()
Static Public Member Functions inherited from ck_tile::tile_window_base< tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >, BottomTensorView_, WindowLengths_ >
static constexpr CK_TILE_DEVICE index_t	get_num_of_dimension ()

Public Attributes
array< typename Base::BottomTensorCoord, traits::NumAccess_NonLinear >	cached_coords_
array< typename Base::WindowAdaptorCoord, traits::NumAccess_NonLinear >	cached_window_adaptor_coords_
array< bool, Base::Traits::NumAccess >	cached_flags_
Public Attributes inherited from ck_tile::tile_window_with_tile_dstr_base< tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >, BottomTensorView_, WindowLengths_, StaticTileDistribution_ >
TileDstr	tile_dstr_
Public Attributes inherited from ck_tile::tile_window_base< tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >, BottomTensorView_, WindowLengths_ >
BottomTensorIndex	window_origin_
WindowLengths	window_lengths_
BottomTensorView	bottom_tensor_view_

Static Public Attributes
static constexpr auto	I0 = number<0>{}
static constexpr auto	I1 = number<1>{}
static constexpr index_t	NumAccess = Base::Traits::NumAccess
static constexpr index_t	NumAccess_NonLinear = traits::NumAccess_NonLinear
Static Public Attributes inherited from ck_tile::tile_window_with_tile_dstr_base< tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >, BottomTensorView_, WindowLengths_, StaticTileDistribution_ >
static constexpr index_t	NDimWindowAdaptorTop
static constexpr index_t	NDimP
static constexpr index_t	NDimY
Static Public Attributes inherited from ck_tile::tile_window_base< tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >, BottomTensorView_, WindowLengths_ >
static constexpr index_t	NDimBottomTensor

Member Typedef Documentation

◆ AccessHistogram_NonLinear

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

using ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::AccessHistogram_NonLinear = typename traits::AccessHistogram_NonLinear

◆ AccessMap_NonLinear

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

using ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::AccessMap_NonLinear = typename traits::AccessMap_NonLinear

◆ AccessPrefixSum_NonLinear

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

using ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::AccessPrefixSum_NonLinear = typename traits::AccessPrefixSum_NonLinear

◆ Base

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

using ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::Base = tile_window_with_tile_dstr_base<tile_window_linear<BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_>, BottomTensorView_, WindowLengths_, StaticTileDistribution_>

◆ LinearBottomDims

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

using ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::LinearBottomDims = remove_cvref_t<LinearBottomDims_>

Constructor & Destructor Documentation

◆ tile_window_linear() [1/2]

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

constexpr CK_TILE_DEVICE ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::tile_window_linear ( )

constexprdefault

◆ tile_window_linear() [2/2]

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

constexpr CK_TILE_DEVICE ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::tile_window_linear ( const typename Base::BottomTensorView &  bottom_tensor_view, const typename Base::WindowLengths &  window_lengths, const typename Base::BottomTensorIndex &  window_origin, const typename Base::TileDstr &  tile_distribution  )

inlineconstexpr

Member Function Documentation

◆ async_load()

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

template<typename LdsTileWindow_ , index_t i_access = -1, bool oob_conditional_check = true>

CK_TILE_DEVICE auto ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::async_load ( LdsTileWindow_ &&  lds_tile, number< i_access >  = {}, bool_constant< oob_conditional_check >  = {}  ) const

inline

◆ async_load_raw()

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

template<typename LdsTileWindow_ , index_t i_access = -1, bool oob_conditional_check = true, bool pre_nop = false>

CK_TILE_DEVICE auto ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::async_load_raw ( LdsTileWindow_ &&  lds_tile, number< i_access >  = {}, bool_constant< oob_conditional_check >  = {}, bool_constant< pre_nop >  = {}  ) const

inline

◆ get_bottom_linear_coordinate()

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

template<index_t i_access>

static constexpr CK_TILE_DEVICE auto ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::get_bottom_linear_coordinate ( number< i_access >  )

inlinestaticconstexpr

◆ get_bottom_linear_offset()

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

template<index_t i_access>

static constexpr CK_TILE_DEVICE index_t ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::get_bottom_linear_offset ( number< i_access >  )

inlinestaticconstexpr

◆ load() [1/2]

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

template<typename DstTile , index_t i_access = -1, bool oob_conditional_check = true>

CK_TILE_DEVICE auto ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::load ( DstTile &  dst_tensor, number< i_access >  = {}, bool_constant< oob_conditional_check >  = {}  ) const

inline

◆ load() [2/2]

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

template<index_t i_access = -1, bool oob_conditional_check = true>

CK_TILE_DEVICE auto ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::load ( number< i_access >  = {}, bool_constant< oob_conditional_check >  = {}  ) const

inline

◆ load_raw()

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

template<typename DstTile , index_t i_access = -1, bool oob_conditional_check = true, bool pre_nop = false>

CK_TILE_DEVICE void ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::load_raw ( DstTile &  dst_tensor, number< i_access >  = {}, bool_constant< oob_conditional_check >  = {}, bool_constant< pre_nop >  = {}  ) const

inline

◆ load_transpose()

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

template<typename Policy , index_t i_access_unsupport_ = -1, bool oob_conditional_check = true>

CK_TILE_DEVICE auto ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::load_transpose ( ) const

inline

◆ load_transpose_linear()

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

template<typename Policy , typename DistributedTensor , index_t i_access = -1, bool oob_conditional_check = true>

CK_TILE_DEVICE auto ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::load_transpose_linear ( DistributedTensor &  dst_tensor, number< i_access >  = {}, bool_constant< oob_conditional_check >  = {}  ) const

inline

◆ move_extended()

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

CK_TILE_DEVICE void ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::move_extended ( const typename Base::BottomTensorIndex &  step )

inline

◆ set_window_origin_extended()

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

CK_TILE_DEVICE void ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::set_window_origin_extended ( const typename Base::BottomTensorIndex &  )

inline

◆ store()

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

template<index_t i_access = -1, bool oob_conditional_check = true>

CK_TILE_DEVICE void ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::store ( const static_distributed_tensor< typename Base::DataType, typename Base::TileDstr > &  dstr_tensor, number< i_access >  = {}, bool_constant< oob_conditional_check >  = {}  ) const

inline

◆ store_raw()

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

template<index_t i_access = -1>

CK_TILE_DEVICE void ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::store_raw ( const static_distributed_tensor< typename Base::DataType, typename Base::TileDstr > &  dstr_tensor, number< i_access >  = {}  ) const

inline

◆ update()

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

template<index_t i_access = -1, bool oob_conditional_check = true>

CK_TILE_DEVICE void ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::update ( const static_distributed_tensor< typename Base::DataType, typename Base::TileDstr > &  dstr_tensor, number< i_access >  = {}, bool_constant< oob_conditional_check >  = {}  ) const

inline

◆ update_raw()

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

template<index_t i_access = -1, bool oob_conditional_check = true, bool pre_nop = false>

CK_TILE_DEVICE void ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::update_raw ( const static_distributed_tensor< typename Base::DataType, typename Base::TileDstr > &  dstr_tensor, number< i_access >  = {}, bool_constant< oob_conditional_check >  = {}, bool_constant< pre_nop >  = {}  ) const

inline

Member Data Documentation

◆ cached_coords_

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

array<typename Base::BottomTensorCoord, traits::NumAccess_NonLinear> ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::cached_coords_

◆ cached_flags_

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

array<bool, Base::Traits::NumAccess> ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::cached_flags_

◆ cached_window_adaptor_coords_

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

array<typename Base::WindowAdaptorCoord, traits::NumAccess_NonLinear> ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::cached_window_adaptor_coords_

◆ I0

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

constexpr auto ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::I0 = number<0>{}

staticconstexpr

◆ I1

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

constexpr auto ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::I1 = number<1>{}

staticconstexpr

◆ NumAccess

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

constexpr index_t ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::NumAccess = Base::Traits::NumAccess

staticconstexpr

◆ NumAccess_NonLinear

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >

constexpr index_t ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::NumAccess_NonLinear = traits::NumAccess_NonLinear

staticconstexpr

---

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

• /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/core/tensor/tile_window_linear.hpp