layout_utils.hpp File Reference

layout_utils.hpp File Reference#

Composable Kernel: 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__ __device__ constexpr auto	make_layout (const Shape &shape, const Strides &strides)
	Make layout function.
template<typename Shape>
__host__ __device__ constexpr auto	make_layout (const Shape &shape)
	Make layout function with packed strides (column-major).
template<index_t idx, typename... Dims>
__host__ __device__ constexpr auto	get (const Tuple< Dims... > &tuple)
	Get element from tuple (Shape/Strides/Idxs).
template<index_t idx, typename Shape, typename UnrolledDesc>
__host__ __device__ constexpr auto	get (const Layout< Shape, UnrolledDesc > &layout)
	Get sub layout.
template<index_t Idx, index_t... Idxs, typename T>
__host__ __device__ constexpr auto	get (const T &elem)
	Hierarchical get.
template<index_t idx, typename Shape, typename UnrolledDescriptorType>
__host__ __device__ constexpr auto	size (const Layout< Shape, UnrolledDescriptorType > &layout)
	Length get (product if tuple).
template<typename... ShapeDims>
__host__ __device__ constexpr auto	size (const Tuple< ShapeDims... > &shape)
	Shape size (product of dims).
template<typename Shape, typename UnrolledDescriptorType>
__host__ __device__ constexpr auto	size (const Layout< Shape, UnrolledDescriptorType > &layout)
	Layout size (product of dims).
template<index_t idx, typename... Ts>
__host__ __device__ constexpr auto	size (const Tuple< Ts... > &tuple)
	Length get from tuple (product if tuple).
template<index_t Idx, index_t... Idxs, typename T>
__host__ __device__ constexpr auto	size (const T &elem)
	Hierarchical size.
template<typename Shape, typename UnrolledDescriptorType>
__host__ __device__ constexpr auto	rank (const Layout< Shape, UnrolledDescriptorType > &layout)
	Get layout rank (num elements in shape).
template<typename... Dims>
__host__ __device__ constexpr auto	rank (const Tuple< Dims... > &tuple)
	Get tuple rank (num elements in tuple). Return 1 if scalar passed.
template<index_t... Idxs, typename T>
__host__ __device__ constexpr auto	rank (const T &elem)
	Hierarchical rank.
template<typename Shape, typename UnrolledDescriptorType>
__host__ __device__ constexpr auto	depth (const Layout< Shape, UnrolledDescriptorType > &layout)
	Get depth of the layout shape (return 0 if scalar).
template<typename... Dims>
__host__ __device__ constexpr auto	depth (const Tuple< Dims... > &tuple)
	Get depth of the tuple. (return 0 if scalar).
template<index_t... Idxs, typename T>
__host__ __device__ constexpr auto	depth (const T &elem)
	Hierarchical depth.
template<typename LayoutType>
__host__ __device__ constexpr const auto &	shape (const LayoutType &layout)
	Get Layout shape.
template<typename Shape, typename UnrolledDesc, typename TileLengths>
__host__ __device__ constexpr auto	pad (const Layout< Shape, UnrolledDesc > &layout, const TileLengths &tile_lengths)
	Pad layout shapes to be adjusted to tile lengths.
template<index_t Idx, typename Shape, typename UnrolledDesc, typename NewLengths, typename NewIdxs>
__host__ __device__ constexpr auto	unmerge (const Layout< Shape, UnrolledDesc > &layout, const NewLengths &new_lengths, const NewIdxs &new_indexes)
	Unmerge selected dim in layout.

Function Documentation

◆ depth() [1/3]

template<typename Shape, typename UnrolledDescriptorType>

__host__ __device__ constexpr 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__ __device__ constexpr 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__ __device__ constexpr 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__ __device__ constexpr 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__ __device__ constexpr 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__ __device__ constexpr 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__ __device__ constexpr 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__ __device__ constexpr 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__ __device__ constexpr 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__ __device__ constexpr auto rank ( 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<index_t... Idxs, typename T>

__host__ __device__ constexpr auto rank ( const T & elem )

constexpr

Hierarchical rank.

Template Parameters

Idxs	Indexes to lookup.

Parameters

elem	Element to lookup.

Returns: Requsted rank.

◆ rank() [3/3]

template<typename... Dims>

__host__ __device__ constexpr auto rank ( 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.

◆ shape()

template<typename LayoutType>

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

constexpr

Get Layout shape.

Parameters

layout Layout to get shape from.

Returns: Requsted shape.

◆ size() [1/5]

template<typename Shape, typename UnrolledDescriptorType>

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

constexpr

Layout size (product of dims).

Parameters

layout Layout to calculate shape size.

Returns: Requsted size.

◆ size() [2/5]

template<index_t idx, typename Shape, typename UnrolledDescriptorType>

__host__ __device__ constexpr 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() [3/5]

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

__host__ __device__ constexpr 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__ __device__ constexpr 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__ __device__ constexpr 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__ __device__ constexpr auto unmerge	(	const Layout< Shape, UnrolledDesc > &	layout,
		const NewLengths &	new_lengths,
		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.

layout_utils.hpp File Reference

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()