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

tile_window_linear&lt; BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ &gt; Struct Template Reference#

Composable Kernel: ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ > Struct Template Reference
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>

Classes

struct  traits
 

Public Types

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

Public Member Functions

constexpr CK_TILE_DEVICE tile_window_linear ()=default
 
constexpr CK_TILE_DEVICE tile_window_linear (const BottomTensorView &bottom_tensor_view, const WindowLengths &window_lengths, const BottomTensorIndex &window_origin, const TileDstr &tile_distribution)
 
constexpr CK_TILE_DEVICE auto get_window_lengths () const
 
constexpr CK_TILE_DEVICE auto get_tile_distribution () const
 
constexpr CK_TILE_DEVICE auto get_bottom_tensor_view () const
 
constexpr CK_TILE_DEVICE auto get_window_origin () const
 
constexpr CK_TILE_DEVICE void set_bottom_tensor_view_data_ptr (typename BottomTensorView::DataType *data)
 
template<typename ATopIndex >
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
 
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<index_t i_access = -1, bool oob_conditional_check = true>
CK_TILE_DEVICE void store (const static_distributed_tensor< DataType, 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< DataType, 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< DataType, 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< DataType, TileDstr > &dstr_tensor, number< i_access >={}, bool_constant< oob_conditional_check >={}, bool_constant< pre_nop >={}) const
 
CK_TILE_DEVICE void move (const BottomTensorIndex &step)
 
CK_TILE_DEVICE void set_window_origin (const BottomTensorIndex &new_window_origin)
 
CK_TILE_HOST_DEVICE void init_raw ()
 

Static Public Member Functions

static constexpr CK_TILE_DEVICE index_t get_num_of_dimension ()
 
static constexpr CK_TILE_DEVICE bool has_static_tile_distribution ()
 
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 >)
 

Public Attributes

BottomTensorView bottom_tensor_view_
 
WindowLengths window_lengths_
 
BottomTensorIndex window_origin_
 
TileDstr tile_dstr_
 
array< BottomTensorCoord, traits::NumAccess_NonLinear > cached_coords_
 
array< bool, traits::NumAccess > cached_flags_
 

Static Public Attributes

static constexpr index_t NDimWindowAdaptorTop = WindowAdaptor::get_num_of_top_dimension()
 
static constexpr index_t NDimBottomTensor = BottomTensorDesc::get_num_of_dimension()
 
static constexpr index_t NDimP = TileDstr::get_num_of_dimension_p()
 
static constexpr index_t NDimY = TileDstr::get_num_of_dimension_y()
 
static constexpr auto I0 = number<0>{}
 
static constexpr auto I1 = number<1>{}
 
static constexpr index_t NumAccess = traits::NumAccess
 
static constexpr index_t NumAccess_NonLinear = traits::NumAccess_NonLinear
 

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

◆ AdaptorTopIndex

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
using ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::AdaptorTopIndex = array<index_t, NDimWindowAdaptorTop>

◆ BottomTensorCoord

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
using ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::BottomTensorCoord = decltype(make_tensor_coordinate(BottomTensorDesc{}, BottomTensorIndex{}))

◆ BottomTensorDesc

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
using ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::BottomTensorDesc = typename BottomTensorView::TensorDesc

◆ BottomTensorIndex

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
using ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::BottomTensorIndex = array<index_t, NDimBottomTensor>

◆ BottomTensorView

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
using ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::BottomTensorView = remove_reference_t<BottomTensorView_>

◆ DataType

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
using ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::DataType = remove_cvref_t<typename BottomTensorView::DataType>

◆ 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_>

◆ TileDstr

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
using ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::TileDstr = remove_cvref_t<StaticTileDistribution_>

◆ WindowAdaptor

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
using ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::WindowAdaptor = typename TileDstr::PsYs2XsAdaptor

◆ WindowAdaptorCoord

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
using ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::WindowAdaptorCoord = decltype(make_tensor_adaptor_coordinate(WindowAdaptor{}, AdaptorTopIndex{}))

◆ WindowLengths

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
using ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::WindowLengths = remove_cvref_t<WindowLengths_>

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 BottomTensorView &  bottom_tensor_view,
const WindowLengths &  window_lengths,
const BottomTensorIndex &  window_origin,
const 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

◆ get_bottom_tensor_view()

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
constexpr CK_TILE_DEVICE auto ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::get_bottom_tensor_view ( ) const
inlineconstexpr

◆ get_num_of_access()

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
constexpr CK_TILE_DEVICE auto ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::get_num_of_access ( ) const
inlineconstexpr

◆ get_num_of_dimension()

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
static constexpr CK_TILE_DEVICE index_t ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::get_num_of_dimension ( )
inlinestaticconstexpr

◆ get_tile_distribution()

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
constexpr CK_TILE_DEVICE auto ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::get_tile_distribution ( ) const
inlineconstexpr

◆ get_window_lengths()

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
constexpr CK_TILE_DEVICE auto ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::get_window_lengths ( ) const
inlineconstexpr

◆ get_window_origin()

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
constexpr CK_TILE_DEVICE auto ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::get_window_origin ( ) const
inlineconstexpr

◆ has_static_tile_distribution()

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
static constexpr CK_TILE_DEVICE bool ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::has_static_tile_distribution ( )
inlinestaticconstexpr

◆ init_raw()

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
CK_TILE_HOST_DEVICE void ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::init_raw ( )
inline

◆ 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

◆ move()

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
CK_TILE_DEVICE void ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::move ( const BottomTensorIndex &  step)
inline

◆ move_window_adaptor_and_bottom_tensor_thread_coordinate()

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
template<typename ATopIndex >
CK_TILE_DEVICE void ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::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
inline

◆ set_bottom_tensor_view_data_ptr()

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
constexpr CK_TILE_DEVICE void ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::set_bottom_tensor_view_data_ptr ( typename BottomTensorView::DataType *  data)
inlineconstexpr

◆ set_window_origin()

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 ( const BottomTensorIndex &  new_window_origin)
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< DataType, 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< DataType, 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< DataType, 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< DataType, TileDstr > &  dstr_tensor,
number< i_access >  = {},
bool_constant< oob_conditional_check >  = {},
bool_constant< pre_nop >  = {} 
) const
inline

Member Data Documentation

◆ bottom_tensor_view_

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
BottomTensorView ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::bottom_tensor_view_

◆ cached_coords_

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
array<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, traits::NumAccess> ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::cached_flags_

◆ 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

◆ NDimBottomTensor

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
constexpr index_t ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::NDimBottomTensor = BottomTensorDesc::get_num_of_dimension()
staticconstexpr

◆ NDimP

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
constexpr index_t ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::NDimP = TileDstr::get_num_of_dimension_p()
staticconstexpr

◆ NDimWindowAdaptorTop

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
constexpr index_t ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::NDimWindowAdaptorTop = WindowAdaptor::get_num_of_top_dimension()
staticconstexpr

◆ NDimY

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
constexpr index_t ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::NDimY = TileDstr::get_num_of_dimension_y()
staticconstexpr

◆ NumAccess

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
constexpr index_t ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::NumAccess = 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

◆ tile_dstr_

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
TileDstr ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::tile_dstr_

◆ window_lengths_

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
WindowLengths ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::window_lengths_

◆ window_origin_

template<typename BottomTensorView_ , typename WindowLengths_ , typename StaticTileDistribution_ , typename LinearBottomDims_ >
BottomTensorIndex ck_tile::tile_window_linear< BottomTensorView_, WindowLengths_, StaticTileDistribution_, LinearBottomDims_ >::window_origin_
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/tensor/tile_window_linear.hpp