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

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

Composable Kernel: 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 = decltype(make_tensor_adaptor_coordinate(WindowAdaptor{}, AdaptorTopIndex{}))

using	BottomTensorCoord = decltype(make_tensor_coordinate(BottomTensorDesc{}, BottomTensorIndex{}))

Public Member Functions
constexpr CK_TILE_DEVICE	tile_scatter_gather ()=default

constexpr CK_TILE_DEVICE	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)

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_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, 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	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 constexpr CK_TILE_DEVICE index_t	get_num_of_dimension ()

static constexpr CK_TILE_DEVICE bool	has_static_tile_distribution ()

static constexpr CK_TILE_DEVICE 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
NumCoord	TBD