tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim > Struct Template Reference

tile_scatter_gather&lt; BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim &gt; Struct Template Reference#

Composable Kernel: ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim > 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_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim > Struct Template Reference

This class provides tile (windowed) view and access to the device memory. More...

#include <tile_scatter_gather.hpp>

Classes

struct  load_store_traits

Public Types

using BottomTensorView = remove_reference_t<BottomTensorView_>
using WindowLengths = remove_cvref_t<WindowLengths_>
using TileDstr = remove_cvref_t<StaticTileDistribution_>
using PageIdxArray = remove_cvref_t<StaticPageIndexArray_>
using ValidArray = remove_cvref_t<StaticValidArray_>
using WindowAdaptor = typename TileDstr::PsYs2XsAdaptor
using BottomTensorDesc = typename BottomTensorView::TensorDesc
using DataType = remove_cvref_t<typename BottomTensorView::DataType>
using AdaptorTopIndex = array<index_t, NDimWindowAdaptorTop>
using BottomTensorIndex = array<index_t, NDimBottomTensor>
using WindowAdaptorCoord
using BottomTensorCoord

Public Member Functions

CK_TILE_DEVICE constexpr tile_scatter_gather ()=default
CK_TILE_DEVICE constexpr tile_scatter_gather (const BottomTensorView &bottom_tensor_view, const WindowLengths &window_lengths, const BottomTensorIndex &window_origin, const TileDstr &tile_distribution, const PageIdxArray &page_idx, const ValidArray &valids)
CK_TILE_DEVICE constexpr auto get_window_lengths () const
CK_TILE_DEVICE constexpr auto get_tile_distribution () const
CK_TILE_DEVICE constexpr auto get_bottom_tensor_view () const
CK_TILE_DEVICE constexpr auto get_window_origin () const
CK_TILE_DEVICE constexpr 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
CK_TILE_DEVICE constexpr auto get_num_of_access () const
template<index_t i_access_unsupport_ = -1, bool oob_conditional_check = true>
CK_TILE_DEVICE auto load (number< i_access_unsupport_ >={}, bool_constant< oob_conditional_check >={}) const
template<typename DistributedTensor, index_t i_access_unsupport_ = -1, bool oob_conditional_check = true>
CK_TILE_DEVICE auto load (DistributedTensor &dst_tensor, number< i_access_unsupport_ >={}, bool_constant< oob_conditional_check >={}) const
template<typename LdsTileWindow_, index_t i_access_unsupport_ = -1, bool oob_conditional_check = true>
CK_TILE_DEVICE auto async_load (LdsTileWindow_ &&lds_tile, number< i_access_unsupport_ >={}, bool_constant< oob_conditional_check >={}) const
template<typename LdsTileWindow_, index_t i_access_unsupport_ = -1, bool oob_conditional_check = true, bool pre_nop = false>
CK_TILE_DEVICE auto async_load_raw (LdsTileWindow_ &&lds_tile, number< i_access_unsupport_ >={}, bool_constant< oob_conditional_check >={}, bool_constant< pre_nop >={}) const
template<index_t i_access_unsupport_ = -1, bool oob_conditional_check = true>
CK_TILE_DEVICE void update (const static_distributed_tensor< DataType, TileDstr > &dstr_tensor, number< i_access_unsupport_ >={}, bool_constant< oob_conditional_check >={}) const
template<index_t i_access_unsupport_ = -1, bool oob_conditional_check = true>
CK_TILE_DEVICE void store (const static_distributed_tensor< DataType, TileDstr > &dstr_tensor, number< i_access_unsupport_ >={}, bool_constant< oob_conditional_check >={}) const
CK_TILE_DEVICE void move (const BottomTensorIndex &step)
CK_TILE_DEVICE void update_page_idx (const PageIdxArray &new_idx)
CK_TILE_DEVICE void update_valids (const ValidArray &new_valids)
CK_TILE_DEVICE void update_page_idx_and_valids (const PageIdxArray &new_idx, const ValidArray &new_valids)
CK_TILE_DEVICE void set_window_origin (const BottomTensorIndex &new_window_origin)
CK_TILE_HOST_DEVICE void init_raw ()

Static Public Member Functions

static CK_TILE_DEVICE constexpr index_t get_num_of_dimension ()
static CK_TILE_DEVICE constexpr bool has_static_tile_distribution ()
static CK_TILE_DEVICE constexpr auto get_window_adaptor_ys_safe_vector_length_strides ()

Public Attributes

BottomTensorView bottom_tensor_view_
WindowLengths window_lengths_
BottomTensorIndex window_origin_
TileDstr tile_dstr_
PageIdxArray page_idx_
ValidArray valids_
array< tuple< WindowAdaptorCoord, BottomTensorCoord >, NumCoord > pre_computed_coords_

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 NumAccessPerCoord = load_store_traits::NumAccess / NumCoord

Detailed Description

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
struct ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >

This class provides tile (windowed) view and access to the device memory.

Note
This tile window does not support single issue you need to use tile_window_linear structure for this purpose
Template Parameters
BottomTensorView_Class describing & holding device tensor memory.
WindowLengths_Spatial sizes of windowed view on tensor.
StaticTileDistribution_Thread distribution (mapping) into Tile dimensions
NumCoordTBD

Member Typedef Documentation

◆ AdaptorTopIndex

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
using ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::AdaptorTopIndex = array<index_t, NDimWindowAdaptorTop>

◆ BottomTensorCoord

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
using ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::BottomTensorCoord
Initial value:
decltype(make_tensor_coordinate(BottomTensorDesc{}, BottomTensorIndex{}))
ck_tile::make_tensor_coordinate
CK_TILE_HOST_DEVICE constexpr auto make_tensor_coordinate(const TensorDesc &tensor_desc, const TopIndex &idx_top)
Definition tensor_coordinate.hpp:60
ck_tile::tile_scatter_gather::BottomTensorIndex
array< index_t, NDimBottomTensor > BottomTensorIndex
Definition tile_scatter_gather.hpp:73
ck_tile::tile_scatter_gather::BottomTensorDesc
typename BottomTensorView::TensorDesc BottomTensorDesc
Definition tile_scatter_gather.hpp:49

◆ BottomTensorDesc

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
using ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::BottomTensorDesc = typename BottomTensorView::TensorDesc

◆ BottomTensorIndex

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
using ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::BottomTensorIndex = array<index_t, NDimBottomTensor>

◆ BottomTensorView

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
using ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::BottomTensorView = remove_reference_t<BottomTensorView_>

◆ DataType

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
using ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::DataType = remove_cvref_t<typename BottomTensorView::DataType>

◆ PageIdxArray

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
using ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::PageIdxArray = remove_cvref_t<StaticPageIndexArray_>

◆ TileDstr

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
using ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::TileDstr = remove_cvref_t<StaticTileDistribution_>

◆ ValidArray

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
using ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::ValidArray = remove_cvref_t<StaticValidArray_>

◆ WindowAdaptor

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
using ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::WindowAdaptor = typename TileDstr::PsYs2XsAdaptor

◆ WindowAdaptorCoord

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
using ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::WindowAdaptorCoord
Initial value:
decltype(make_tensor_adaptor_coordinate(WindowAdaptor{}, AdaptorTopIndex{}))
ck_tile::make_tensor_adaptor_coordinate
CK_TILE_HOST_DEVICE constexpr auto make_tensor_adaptor_coordinate(const Adaptor &adaptor, const TopIndex &idx_top)
Definition tensor_adaptor_coordinate.hpp:55
ck_tile::tile_scatter_gather::WindowAdaptor
typename TileDstr::PsYs2XsAdaptor WindowAdaptor
Definition tile_scatter_gather.hpp:48
ck_tile::tile_scatter_gather::AdaptorTopIndex
array< index_t, NDimWindowAdaptorTop > AdaptorTopIndex
Definition tile_scatter_gather.hpp:72

◆ WindowLengths

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
using ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::WindowLengths = remove_cvref_t<WindowLengths_>

Constructor & Destructor Documentation

◆ tile_scatter_gather() [1/2]

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
CK_TILE_DEVICE constexpr ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::tile_scatter_gather ( )
constexprdefault

◆ tile_scatter_gather() [2/2]

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
CK_TILE_DEVICE constexpr ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::tile_scatter_gather ( const BottomTensorView & bottom_tensor_view,
const WindowLengths & window_lengths,
const BottomTensorIndex & window_origin,
const TileDstr & tile_distribution,
const PageIdxArray & page_idx,
const ValidArray & valids )
inlineconstexpr

Member Function Documentation

◆ async_load()

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
template<typename LdsTileWindow_, index_t i_access_unsupport_ = -1, bool oob_conditional_check = true>
CK_TILE_DEVICE auto ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::async_load ( LdsTileWindow_ && lds_tile,
number< i_access_unsupport_ > = {},
bool_constant< oob_conditional_check > = {} ) const
inline

TODO: use structure binding (to be captured later) if compiled in C++20

◆ async_load_raw()

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
template<typename LdsTileWindow_, index_t i_access_unsupport_ = -1, bool oob_conditional_check = true, bool pre_nop = false>
CK_TILE_DEVICE auto ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::async_load_raw ( LdsTileWindow_ && lds_tile,
number< i_access_unsupport_ > = {},
bool_constant< oob_conditional_check > = {},
bool_constant< pre_nop > = {} ) const
inline

TODO: use structure binding (to be captured later) if compiled in C++20

◆ get_bottom_tensor_view()

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
CK_TILE_DEVICE constexpr auto ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::get_bottom_tensor_view ( ) const
inlineconstexpr

◆ get_num_of_access()

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
CK_TILE_DEVICE constexpr auto ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::get_num_of_access ( ) const
inlineconstexpr

◆ get_num_of_dimension()

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
CK_TILE_DEVICE constexpr index_t ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::get_num_of_dimension ( )
inlinestaticconstexpr

◆ get_tile_distribution()

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
CK_TILE_DEVICE constexpr auto ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::get_tile_distribution ( ) const
inlineconstexpr

◆ get_window_adaptor_ys_safe_vector_length_strides()

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
CK_TILE_DEVICE constexpr auto ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::get_window_adaptor_ys_safe_vector_length_strides ( )
inlinestaticconstexpr

◆ get_window_lengths()

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
CK_TILE_DEVICE constexpr auto ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::get_window_lengths ( ) const
inlineconstexpr

◆ get_window_origin()

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
CK_TILE_DEVICE constexpr auto ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::get_window_origin ( ) const
inlineconstexpr

◆ has_static_tile_distribution()

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
CK_TILE_DEVICE constexpr bool ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::has_static_tile_distribution ( )
inlinestaticconstexpr

◆ init_raw()

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
CK_TILE_HOST_DEVICE void ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::init_raw ( )
inline

◆ load() [1/2]

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
template<typename DistributedTensor, index_t i_access_unsupport_ = -1, bool oob_conditional_check = true>
CK_TILE_DEVICE auto ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::load ( DistributedTensor & dst_tensor,
number< i_access_unsupport_ > = {},
bool_constant< oob_conditional_check > = {} ) const
inline

TODO: use structure binding (to be captured later) if compiled in C++20

◆ load() [2/2]

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
template<index_t i_access_unsupport_ = -1, bool oob_conditional_check = true>
CK_TILE_DEVICE auto ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::load ( number< i_access_unsupport_ > = {},
bool_constant< oob_conditional_check > = {} ) const
inline

◆ move()

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
CK_TILE_DEVICE void ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::move ( const BottomTensorIndex & step)
inline

◆ move_window_adaptor_and_bottom_tensor_thread_coordinate()

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
template<typename ATopIndex>
CK_TILE_DEVICE void ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::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 StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
CK_TILE_DEVICE constexpr void ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::set_bottom_tensor_view_data_ptr ( typename BottomTensorView::DataType * data)
inlineconstexpr

◆ set_window_origin()

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
CK_TILE_DEVICE void ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::set_window_origin ( const BottomTensorIndex & new_window_origin)
inline

◆ store()

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
template<index_t i_access_unsupport_ = -1, bool oob_conditional_check = true>
CK_TILE_DEVICE void ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::store ( const static_distributed_tensor< DataType, TileDstr > & dstr_tensor,
number< i_access_unsupport_ > = {},
bool_constant< oob_conditional_check > = {} ) const
inline

◆ update()

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
template<index_t i_access_unsupport_ = -1, bool oob_conditional_check = true>
CK_TILE_DEVICE void ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::update ( const static_distributed_tensor< DataType, TileDstr > & dstr_tensor,
number< i_access_unsupport_ > = {},
bool_constant< oob_conditional_check > = {} ) const
inline

◆ update_page_idx()

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
CK_TILE_DEVICE void ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::update_page_idx ( const PageIdxArray & new_idx)
inline

◆ update_page_idx_and_valids()

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
CK_TILE_DEVICE void ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::update_page_idx_and_valids ( const PageIdxArray & new_idx,
const ValidArray & new_valids )
inline

◆ update_valids()

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
CK_TILE_DEVICE void ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::update_valids ( const ValidArray & new_valids)
inline

Member Data Documentation

◆ bottom_tensor_view_

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
BottomTensorView ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::bottom_tensor_view_

◆ I0

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
auto ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::I0 = number<0>{}
staticconstexpr

◆ I1

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
auto ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::I1 = number<1>{}
staticconstexpr

◆ NDimBottomTensor

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
index_t ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::NDimBottomTensor = BottomTensorDesc::get_num_of_dimension()
staticconstexpr

◆ NDimP

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
index_t ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::NDimP = TileDstr::get_num_of_dimension_p()
staticconstexpr

◆ NDimWindowAdaptorTop

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
index_t ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::NDimWindowAdaptorTop = WindowAdaptor::get_num_of_top_dimension()
staticconstexpr

◆ NDimY

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
index_t ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::NDimY = TileDstr::get_num_of_dimension_y()
staticconstexpr

◆ NumAccessPerCoord

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
index_t ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::NumAccessPerCoord = load_store_traits::NumAccess / NumCoord
staticconstexpr

◆ page_idx_

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
PageIdxArray ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::page_idx_

◆ pre_computed_coords_

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
array<tuple<WindowAdaptorCoord, BottomTensorCoord>, NumCoord> ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::pre_computed_coords_

◆ tile_dstr_

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
TileDstr ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::tile_dstr_

◆ valids_

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
ValidArray ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::valids_

◆ window_lengths_

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
WindowLengths ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::window_lengths_

◆ window_origin_

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, typename StaticPageIndexArray_, typename StaticValidArray_, index_t HsGatherDim = 0, index_t NumCoord = 1, index_t YsGatherDim = 0>
BottomTensorIndex ck_tile::tile_scatter_gather< BottomTensorView_, WindowLengths_, StaticTileDistribution_, StaticPageIndexArray_, StaticValidArray_, HsGatherDim, NumCoord, YsGatherDim >::window_origin_
The documentation for this struct was generated from the following file: