/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/wrapper/utils/layout_utils.hpp File Reference

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/wrapper/utils/layout_utils.hpp File Reference#

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/wrapper/utils/layout_utils.hpp File Reference

#include "ck/ck.hpp"
#include "ck/utility/number.hpp"
#include "ck/utility/tuple.hpp"
#include "ck/utility/tuple_helper.hpp"
#include "ck/utility/sequence.hpp"
#include "ck/utility/sequence_helper.hpp"
#include "ck/utility/is_detected.hpp"
#include "ck/tensor_description/tensor_descriptor.hpp"
#include "ck/tensor_description/tensor_descriptor_helper.hpp"
#include "ck/tensor_description/multi_index_transform_helper.hpp"
#include "ck/tensor_operation/gpu/device/matrix_padder.hpp"

Go to the source code of this file.

Functions
template<typename Shape , typename Strides >
__host__ constexpr __device__ auto	make_layout (const Shape &shape, const Strides &strides)
	Make layout function. More...

template<typename Shape >
__host__ constexpr __device__ auto	make_layout (const Shape &shape)
	Make layout function with packed strides (column-major). More...

template<index_t idx, typename... Dims>
__host__ constexpr __device__ auto	get (const Tuple< Dims... > &tuple)
	Get element from tuple (Shape/Strides/Idxs). More...

template<index_t idx, typename Shape , typename UnrolledDesc >
__host__ constexpr __device__ auto	get (const Layout< Shape, UnrolledDesc > &layout)
	Get sub layout. More...

template<index_t Idx, index_t... Idxs, typename T >
__host__ constexpr __device__ auto	get (const T &elem)
	Hierarchical get. More...

template<index_t idx, typename Shape , typename UnrolledDescriptorType >
__host__ constexpr __device__ auto	size (const Layout< Shape, UnrolledDescriptorType > &layout)
	Length get (product if tuple). More...

template<typename... ShapeDims>
__host__ constexpr __device__ auto	size (const Tuple< ShapeDims... > &shape)
	Shape size (product of dims). More...

template<typename Shape , typename UnrolledDescriptorType >
__host__ constexpr __device__ auto	size (const Layout< Shape, UnrolledDescriptorType > &layout)
	Layout size (product of dims). More...

template<index_t idx, typename... Ts>
__host__ constexpr __device__ auto	size (const Tuple< Ts... > &tuple)
	Length get from tuple (product if tuple). More...

template<index_t Idx, index_t... Idxs, typename T >
__host__ constexpr __device__ auto	size (const T &elem)
	Hierarchical size. More...

template<typename Shape , typename UnrolledDescriptorType >
__host__ constexpr __device__ auto	rank ([[maybe_unused]] const Layout< Shape, UnrolledDescriptorType > &layout)
	Get layout rank (num elements in shape). More...

template<typename... Dims>
__host__ constexpr __device__ auto	rank ([[maybe_unused]] const Tuple< Dims... > &tuple)
	Get tuple rank (num elements in tuple). Return 1 if scalar passed. More...

template<index_t... Idxs, typename T >
__host__ constexpr __device__ auto	rank (const T &elem)
	Hierarchical rank. More...

template<typename Shape , typename UnrolledDescriptorType >
__host__ constexpr __device__ auto	depth (const Layout< Shape, UnrolledDescriptorType > &layout)
	Get depth of the layout shape (return 0 if scalar). More...

template<typename... Dims>
__host__ constexpr __device__ auto	depth (const Tuple< Dims... > &tuple)
	Get depth of the tuple. (return 0 if scalar) More...

template<index_t... Idxs, typename T >
__host__ constexpr __device__ auto	depth (const T &elem)
	Hierarchical depth. More...

template<typename LayoutType >
__host__ constexpr __device__ const auto &	shape (const LayoutType &layout)
	Get Layout shape. More...

template<typename Shape , typename UnrolledDesc , typename TileLengths >
__host__ constexpr __device__ auto	pad (const Layout< Shape, UnrolledDesc > &layout, const TileLengths &tile_lengths)
	Pad layout shapes to be adjusted to tile lengths. More...

template<index_t Idx, typename Shape , typename UnrolledDesc , typename NewLengths , typename NewIdxs >
__host__ constexpr __device__ auto	unmerge (const Layout< Shape, UnrolledDesc > &layout, const NewLengths &new_lengths, [[maybe_unused]] const NewIdxs &new_indexes)
	Unmerge selected dim in layout. More...

Function Documentation

◆ depth() [1/3]

template<typename Shape , typename UnrolledDescriptorType >

__host__ constexpr __device__ auto depth ( const Layout< Shape, UnrolledDescriptorType > & layout )

constexpr

Get depth of the layout shape (return 0 if scalar).

Parameters

layout Layout to calculate depth.

Returns: Requsted depth.

◆ depth() [2/3]

template<index_t... Idxs, typename T >

__host__ constexpr __device__ auto depth ( const T & elem )

constexpr

Hierarchical depth.

Template Parameters

Idxs	Indexes to lookup.

Parameters

elem	Element to lookup.

Returns: Requsted depth.

◆ depth() [3/3]

template<typename... Dims>

__host__ constexpr __device__ auto depth ( const Tuple< Dims... > & tuple )

constexpr

Get depth of the tuple. (return 0 if scalar)

Parameters

tuple Tuple to calculate depth.

Returns: Requsted depth.

◆ get() [1/3]

template<index_t idx, typename Shape , typename UnrolledDesc >

__host__ constexpr __device__ auto get ( const Layout< Shape, UnrolledDesc > & layout )

constexpr

Get sub layout.

Template Parameters

idx	Index to lookup.

Parameters

layout Layout to create sub layout.

Returns: Requsted sub layout.

◆ get() [2/3]

template<index_t Idx, index_t... Idxs, typename T >

__host__ constexpr __device__ auto get ( const T & elem )

constexpr

Hierarchical get.

Template Parameters

Idxs	Indexes to lookup.

Parameters

elem	Element to lookup.

Returns: Requsted element.

◆ get() [3/3]

template<index_t idx, typename... Dims>

__host__ constexpr __device__ auto get ( const Tuple< Dims... > & tuple )

constexpr

Get element from tuple (Shape/Strides/Idxs).

Template Parameters

idx	Index to lookup.

Parameters

tuple Tuple to lookup.

Returns: Requsted element.

◆ make_layout() [1/2]

template<typename Shape >

__host__ constexpr __device__ auto make_layout ( const Shape & shape )

constexpr

Make layout function with packed strides (column-major).

Template Parameters

Shape Shape for layout.

Returns: Constructed layout.

◆ make_layout() [2/2]

template<typename Shape , typename Strides >

__host__ constexpr __device__ auto make_layout	(	const Shape &	shape,
		const Strides &	strides
	)

constexpr

Make layout function.

Template Parameters

Shape	Shape for layout.
Strides	Strides for layout.

Returns: Constructed layout.

◆ pad()

template<typename Shape , typename UnrolledDesc , typename TileLengths >

__host__ constexpr __device__ auto pad	(	const Layout< Shape, UnrolledDesc > &	layout,
		const TileLengths &	tile_lengths
	)

constexpr

Pad layout shapes to be adjusted to tile lengths.

Parameters

layout	Layout to pad.
tile_lengths	Tile lengths to align layout shape.

Returns: Padded layout.

◆ rank() [1/3]

template<typename Shape , typename UnrolledDescriptorType >

__host__ constexpr __device__ auto rank ( [[maybe_unused] ] const Layout< Shape, UnrolledDescriptorType > & layout )

constexpr

Get layout rank (num elements in shape).

Parameters

layout Layout to calculate rank.

Returns: Requsted rank.

◆ rank() [2/3]

template<typename... Dims>

__host__ constexpr __device__ auto rank ( [[maybe_unused] ] const Tuple< Dims... > & tuple )

constexpr

Get tuple rank (num elements in tuple). Return 1 if scalar passed.

Parameters

tuple Tuple to calculate rank.

Returns: Requsted rank.

◆ rank() [3/3]

template<index_t... Idxs, typename T >

__host__ constexpr __device__ auto rank ( const T & elem )

constexpr

Hierarchical rank.

Template Parameters

Idxs	Indexes to lookup.

Parameters

elem	Element to lookup.

Returns: Requsted rank.

◆ shape()

template<typename LayoutType >

__host__ constexpr __device__ const auto& shape ( const LayoutType & layout )

constexpr

Get Layout shape.

Parameters

layout Layout to get shape from.

Returns: Requsted shape.

◆ size() [1/5]

template<index_t idx, typename Shape , typename UnrolledDescriptorType >

__host__ constexpr __device__ auto size ( const Layout< Shape, UnrolledDescriptorType > & layout )

constexpr

Length get (product if tuple).

Template Parameters

idx	Index to lookup.

Parameters

layout Layout to get Shape of.

Returns: Requsted length.

◆ size() [2/5]

template<typename Shape , typename UnrolledDescriptorType >

__host__ constexpr __device__ auto size ( const Layout< Shape, UnrolledDescriptorType > & layout )

constexpr

Layout size (product of dims).

Parameters

layout Layout to calculate shape size.

Returns: Requsted size.

◆ size() [3/5]

template<index_t Idx, index_t... Idxs, typename T >

__host__ constexpr __device__ auto size ( const T & elem )

constexpr

Hierarchical size.

Template Parameters

Idx	First index to lookup (to avoid empty Idxs).
Idxs	Next indexes to lookup.

Parameters

elem	Element to lookup.

Returns: Requsted element.

◆ size() [4/5]

template<typename... ShapeDims>

__host__ constexpr __device__ auto size ( const Tuple< ShapeDims... > & shape )

constexpr

Shape size (product of dims).

Parameters

shape Shape to lookup.

Returns: Requsted size.

◆ size() [5/5]

template<index_t idx, typename... Ts>

__host__ constexpr __device__ auto size ( const Tuple< Ts... > & tuple )

constexpr

Length get from tuple (product if tuple).

Template Parameters

idx	Index to lookup.

Parameters

tuple Tuple to lookup.

Returns: Requsted length.

◆ unmerge()

template<index_t Idx, typename Shape , typename UnrolledDesc , typename NewLengths , typename NewIdxs >

__host__ constexpr __device__ auto unmerge	(	const Layout< Shape, UnrolledDesc > &	layout,
		const NewLengths &	new_lengths,
		[[maybe_unused] ] const NewIdxs &	new_indexes
	)

constexpr

Unmerge selected dim in layout.

Template Parameters

Idx	Index to dimension being unmerged.

Parameters

layout	Layout to pad.
new_lengths	Dimensions into which the indicated dimension will be divided.
new_indexes	Indexes to shuffle dims. Dims for unmerged dim should be nested.

Returns: Unmerged layout.

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/wrapper/utils/layout_utils.hpp File Reference

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/wrapper/utils/layout_utils.hpp File Reference#

Functions

Function Documentation

◆ depth() [1/3]

◆ depth() [2/3]

◆ depth() [3/3]

◆ get() [1/3]

◆ get() [2/3]

◆ get() [3/3]

◆ make_layout() [1/2]

◆ make_layout() [2/2]

◆ pad()

◆ rank() [1/3]

◆ rank() [2/3]

◆ rank() [3/3]

◆ shape()

◆ size() [1/5]

◆ size() [2/5]

◆ size() [3/5]

◆ size() [4/5]

◆ size() [5/5]

◆ unmerge()