tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord > Struct Template Reference

tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord > Struct Template Reference#

Composable Kernel: ck_tile::tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord > Struct Template Reference

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

#include <tile_window.hpp>

Inheritance diagram for ck_tile::tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord >:

Public Types
using	Base
Public Types inherited from ck_tile::tile_window_with_tile_dstr_base< tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord >, BottomTensorView_, WindowLengths_, StaticTileDistribution_ >
using	TileDstr
using	TileWindowBase
using	WindowAdaptor
using	AdaptorTopIndex
using	WindowAdaptorCoord
using	BottomTensorCoord
Public Types inherited from ck_tile::tile_window_base< TileWindowType_, 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
CK_TILE_DEVICE constexpr	tile_window_with_static_distribution ()=default
CK_TILE_DEVICE constexpr	tile_window_with_static_distribution (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_unsupport_ = -1, bool oob_conditional_check = true>
CK_TILE_DEVICE auto	load (number< i_access_unsupport_ >={}, bool_constant< oob_conditional_check >={}) const
template<typename TileWindow_, typename ElementWise_, index_t i_access_unsupport_ = -1, bool oob_conditional_check = true>
CK_TILE_DEVICE auto	load (const TileWindow_ &tile_window, ElementWise_ elementwise, number< i_access_unsupport_ >={}, bool_constant< oob_conditional_check >={}) const
	Load tile with elementwise function.
template<typename DistributedTensor, typename TileWindow_, typename ElementWise_, index_t i_access_unsupport_ = -1, bool oob_conditional_check = true>
CK_TILE_DEVICE auto	load (DistributedTensor &dst_tensor, const TileWindow_ &tile_window, ElementWise_ elementwise, 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 DstTile, index_t i_access_unsupport_ = -1, bool oob_conditional_check = true, bool pre_nop = false>
CK_TILE_DEVICE void	load_raw (DstTile &dst_tensor, number< i_access_unsupport_ >={}, bool_constant< oob_conditional_check >={}, bool_constant< pre_nop >={}) 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<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 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_unsupport_ = -1, bool oob_conditional_check = true>
CK_TILE_DEVICE auto	load_transpose (DistributedTensor &dst_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< typename Base::DataType, typename Base::TileDstr > &dstr_tensor, number< i_access_unsupport_ >={}, bool_constant< oob_conditional_check >={}) const
template<index_t i_access_unsupport_ = -1>
CK_TILE_DEVICE void	store_raw (const static_distributed_tensor< typename Base::DataType, typename Base::TileDstr > &dstr_tensor, number< i_access_unsupport_ >={}) const
template<index_t i_access_unsupport_ = -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_unsupport_ >={}, bool_constant< oob_conditional_check >={}) const
template<index_t i_access_unsupport_ = -1, bool oob_conditional_check = true, bool pre_nop>
CK_TILE_DEVICE void	update_raw (const static_distributed_tensor< typename Base::DataType, typename Base::TileDstr > &dstr_tensor, number< i_access_unsupport_ >={}, 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_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord >, BottomTensorView_, WindowLengths_, StaticTileDistribution_ >
CK_TILE_DEVICE constexpr 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
CK_TILE_DEVICE constexpr auto	get_num_of_access () const
Public Member Functions inherited from ck_tile::tile_window_base< TileWindowType_, BottomTensorView_, WindowLengths_ >
CK_TILE_DEVICE constexpr auto	get_window_origin () const
CK_TILE_DEVICE constexpr auto	get_window_lengths () const
CK_TILE_DEVICE constexpr 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 &)
CK_TILE_DEVICE constexpr 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 &)

Public Attributes
array< tuple< typename Base::WindowAdaptorCoord, typename Base::BottomTensorCoord >, NumCoord >	pre_computed_coords_
Public Attributes inherited from ck_tile::tile_window_with_tile_dstr_base< tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord >, BottomTensorView_, WindowLengths_, StaticTileDistribution_ >
TileDstr	tile_dstr_
Public Attributes inherited from ck_tile::tile_window_base< TileWindowType_, 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	NumAccessPerCoord = Base::Traits::NumAccess / NumCoord
Static Public Attributes inherited from ck_tile::tile_window_with_tile_dstr_base< tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord >, 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< TileWindowType_, BottomTensorView_, WindowLengths_ >
static constexpr index_t	NDimBottomTensor = BottomTensorDesc::get_num_of_dimension()

Additional Inherited Members
Static Public Member Functions inherited from ck_tile::tile_window_with_tile_dstr_base< tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord >, BottomTensorView_, WindowLengths_, StaticTileDistribution_ >
static CK_TILE_DEVICE constexpr bool	has_static_tile_distribution ()
static CK_TILE_DEVICE constexpr auto	get_window_adaptor_ys_safe_vector_length_strides ()
Static Public Member Functions inherited from ck_tile::tile_window_base< TileWindowType_, BottomTensorView_, WindowLengths_ >
static CK_TILE_DEVICE constexpr index_t	get_num_of_dimension ()

Detailed Description

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, index_t NumCoord>
struct ck_tile::tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord >

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

Member Typedef Documentation

◆ Base

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, index_t NumCoord>

using ck_tile::tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord >::Base

Initial value:

 tile_window_with_tile_dstr_base<
        tile_window_with_static_distribution<BottomTensorView_,
                                             WindowLengths_,
                                             StaticTileDistribution_,
                                             NumCoord>,
        BottomTensorView_,
        WindowLengths_,
        StaticTileDistribution_>

Constructor & Destructor Documentation

◆ tile_window_with_static_distribution() [1/2]

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, index_t NumCoord>

CK_TILE_DEVICE constexpr ck_tile::tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord >::tile_window_with_static_distribution ( )

constexprdefault

◆ tile_window_with_static_distribution() [2/2]

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, index_t NumCoord>

CK_TILE_DEVICE constexpr ck_tile::tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord >::tile_window_with_static_distribution	(	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_, index_t NumCoord>

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

CK_TILE_DEVICE auto ck_tile::tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord >::async_load	(	LdsTileWindow_ &&	lds_tile,
		number< i_access_unsupport_ >	= {},
		bool_constant< oob_conditional_check >	= {} ) const

inline

◆ async_load_raw()

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, index_t NumCoord>

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_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord >::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

◆ load() [1/4]

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, index_t NumCoord>

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

CK_TILE_DEVICE auto ck_tile::tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord >::load	(	const TileWindow_ &	tile_window,
		ElementWise_	elementwise,
		number< i_access_unsupport_ >	= {},
		bool_constant< oob_conditional_check >	= {} ) const

inline

Load tile with elementwise function.

Note: Load tile with elementwise — during value loading, an elementwise function is executed for each A0, A1, … AN. The values A0, A1, … AN are read by the same thread. In this way, we reduce the amount of information loaded into the registers. The same thread, during vectorized reading, accesses the same set of data from A0, A1, A2, … AN.

◆ load() [2/4]

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, index_t NumCoord>

template<typename DistributedTensor, typename TileWindow_, typename ElementWise_, index_t i_access_unsupport_ = -1, bool oob_conditional_check = true>

CK_TILE_DEVICE auto ck_tile::tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord >::load	(	DistributedTensor &	dst_tensor,
		const TileWindow_ &	tile_window,
		ElementWise_	elementwise,
		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() [3/4]

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, index_t NumCoord>

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

CK_TILE_DEVICE auto ck_tile::tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord >::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() [4/4]

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, index_t NumCoord>

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

CK_TILE_DEVICE auto ck_tile::tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord >::load	(	number< i_access_unsupport_ >	= {},
		bool_constant< oob_conditional_check >	= {} ) const

inline

◆ load_raw()

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, index_t NumCoord>

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

CK_TILE_DEVICE void ck_tile::tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord >::load_raw	(	DstTile &	dst_tensor,
		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

◆ load_transpose() [1/2]

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, index_t NumCoord>

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

CK_TILE_DEVICE auto ck_tile::tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord >::load_transpose ( ) const

inline

◆ load_transpose() [2/2]

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, index_t NumCoord>

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

CK_TILE_DEVICE auto ck_tile::tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord >::load_transpose	(	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

◆ move_extended()

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, index_t NumCoord>

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

inline

◆ set_window_origin_extended()

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, index_t NumCoord>

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

inline

◆ store()

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, index_t NumCoord>

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

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

inline

◆ store_raw()

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, index_t NumCoord>

template<index_t i_access_unsupport_ = -1>

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

inline

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

◆ update()

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, index_t NumCoord>

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

CK_TILE_DEVICE void ck_tile::tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord >::update	(	const static_distributed_tensor< typename Base::DataType, typename Base::TileDstr > &	dstr_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

◆ update_raw()

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, index_t NumCoord>

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

CK_TILE_DEVICE void ck_tile::tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord >::update_raw	(	const static_distributed_tensor< typename Base::DataType, typename Base::TileDstr > &	dstr_tensor,
		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

Member Data Documentation

◆ I0

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, index_t NumCoord>

auto ck_tile::tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord >::I0 = number<0>{}

staticconstexpr

◆ I1

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, index_t NumCoord>

auto ck_tile::tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord >::I1 = number<1>{}

staticconstexpr

◆ NumAccessPerCoord

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, index_t NumCoord>

index_t ck_tile::tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord >::NumAccessPerCoord = Base::Traits::NumAccess / NumCoord

staticconstexpr

◆ pre_computed_coords_

template<typename BottomTensorView_, typename WindowLengths_, typename StaticTileDistribution_, index_t NumCoord>

array<tuple<typename Base::WindowAdaptorCoord, typename Base::BottomTensorCoord>, NumCoord> ck_tile::tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord >::pre_computed_coords_

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

tile_window.hpp

tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord > Struct Template Reference

tile_window_with_static_distribution&lt; BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord &gt; Struct Template Reference#

Public Types

Public Member Functions

Public Attributes

Static Public Attributes

Additional Inherited Members

Detailed Description

Member Typedef Documentation

◆ Base

Constructor & Destructor Documentation

◆ tile_window_with_static_distribution() [1/2]

◆ tile_window_with_static_distribution() [2/2]

Member Function Documentation

◆ async_load()

◆ async_load_raw()

◆ load() [1/4]

◆ load() [2/4]

◆ load() [3/4]

◆ load() [4/4]

◆ load_raw()

◆ load_transpose() [1/2]

◆ load_transpose() [2/2]

◆ move_extended()

◆ set_window_origin_extended()

◆ store()

◆ store_raw()

◆ update()

◆ update_raw()

Member Data Documentation

◆ I0

◆ I1

◆ NumAccessPerCoord

◆ pre_computed_coords_

tile_window_with_static_distribution< BottomTensorView_, WindowLengths_, StaticTileDistribution_, NumCoord > Struct Template Reference#