Layout< Shape, UnrolledDescriptorType > Struct Template Reference

Layout< Shape, UnrolledDescriptorType > Struct Template Reference#

Composable Kernel: Layout< Shape, UnrolledDescriptorType > Struct Template Reference

Layout wrapper that performs the tensor descriptor logic. More...

#include <layout.hpp>

Public Types
using	LayoutShape = Shape

using	LayoutUnrolledDescriptorType = UnrolledDescriptorType

using	MergedNestsDescriptorType = remove_cvref_t< decltype(TransformDesc(Shape{}, DefaultIdxsTupleType{}, UnrolledDescriptorType{}))>

Public Member Functions
__host__ constexpr __device__ auto	GetElementSpaceSize () const

__host__ __device__	Layout ()=delete

__host__ constexpr __device__	Layout (const Shape &shape, const UnrolledDescriptorType &unnested_descriptor)
	Layout constructor. More...

template<typename Idxs >
__host__ constexpr __device__ index_t	operator() () const
	Returns real offset to element in runtime. More...

template<typename... Ts>
__host__ __device__ index_t	operator() (const Tuple< Ts... > &Idx) const
	Returns real offset to element in compile time. More...

template<index_t IDim>
__host__ constexpr __device__ auto	GetLength () const
	Length getter (product if tuple). More...

__host__ constexpr __device__ auto	GetLengths () const
	Layout size getter (product of shape). More...

__host__ constexpr __device__ const Shape &	GetShape () const
	Shape getter. More...

__host__ constexpr __device__ auto	GetDefaultLengthsTuple () const
	Get default lengths (tuple filled with Shape length elements). More...

__host__ constexpr __device__ auto	GetDefaultStartIdxs () const
	Get default start idx (tuple filled with 0s of the same size as Shape). More...

__host__ constexpr __device__ const MergedNestsDescriptorType &	GetMergedNestingDescriptor () const
	Get descriptor with all nested dimensions merged. Example, shape: ((2, 2), 2) Descriptor lengths: (4, 2) More...

__host__ constexpr __device__ const Descriptor1dType &	Get1DDescriptor () const
	Get descriptor with all dimensions are merged (1D). Example, shape: ((2, 2), 2) Descriptor lengths: (8) More...

__host__ constexpr __device__ const UnrolledDescriptorType &	GetUnrolledDescriptor () const
	Get unnested descriptor (with unrolled dims) Example, shape: ((2, 2), 2) Descriptor lengths: (2, 2, 2) More...

Static Public Member Functions
template<typename... ShapeDims, typename... IdxDims>
__host__ constexpr static __device__ auto	TransformDesc (const Tuple< ShapeDims... > &shape, const Tuple< IdxDims... > &idxs, const UnrolledDescriptorType &naive_descriptor)
	Transform descriptor to align to passed indexes. More...

Detailed Description

template<typename Shape, typename UnrolledDescriptorType>
struct Layout< Shape, UnrolledDescriptorType >

Layout wrapper that performs the tensor descriptor logic.

Template Parameters

Shape	Tuple of Number<> (for compile-time layout) or index_t (dynamic layout). It is possible to pass nested shapes (e.g. ((4, 2), 2)), nested dimensions are merged.
UnrolledDescriptorType	Tensor descriptor for unnested shape dims.

Member Typedef Documentation

◆ LayoutShape

template<typename Shape , typename UnrolledDescriptorType >

using Layout< Shape, UnrolledDescriptorType >::LayoutShape = Shape

◆ LayoutUnrolledDescriptorType

template<typename Shape , typename UnrolledDescriptorType >

using Layout< Shape, UnrolledDescriptorType >::LayoutUnrolledDescriptorType = UnrolledDescriptorType

◆ MergedNestsDescriptorType

template<typename Shape , typename UnrolledDescriptorType >

using Layout< Shape, UnrolledDescriptorType >::MergedNestsDescriptorType = remove_cvref_t<decltype(TransformDesc( Shape{}, DefaultIdxsTupleType{}, UnrolledDescriptorType{}))>

Constructor & Destructor Documentation

◆ Layout() [1/2]

template<typename Shape , typename UnrolledDescriptorType >

__host__ __device__ Layout< Shape, UnrolledDescriptorType >::Layout ( )

delete

◆ Layout() [2/2]

template<typename Shape , typename UnrolledDescriptorType >

__host__ constexpr __device__ Layout< Shape, UnrolledDescriptorType >::Layout	(	const Shape &	shape,
		const UnrolledDescriptorType &	unnested_descriptor
	)

inlineconstexpr

Layout constructor.

Parameters

shape	Shape for layout.
unnested_descriptor	Descriptor

Member Function Documentation

◆ Get1DDescriptor()

template<typename Shape , typename UnrolledDescriptorType >

__host__ constexpr __device__ const Descriptor1dType& Layout< Shape, UnrolledDescriptorType >::Get1DDescriptor ( ) const

inlineconstexpr

Get descriptor with all dimensions are merged (1D). Example, shape: ((2, 2), 2) Descriptor lengths: (8)

Returns: 1D descriptor.

◆ GetDefaultLengthsTuple()

template<typename Shape , typename UnrolledDescriptorType >

__host__ constexpr __device__ auto Layout< Shape, UnrolledDescriptorType >::GetDefaultLengthsTuple ( ) const

inlineconstexpr

Get default lengths (tuple filled with Shape length elements).

Returns: Default lengths.

◆ GetDefaultStartIdxs()

template<typename Shape , typename UnrolledDescriptorType >

__host__ constexpr __device__ auto Layout< Shape, UnrolledDescriptorType >::GetDefaultStartIdxs ( ) const

inlineconstexpr

Get default start idx (tuple filled with 0s of the same size as Shape).

Returns: Default start idx.

◆ GetElementSpaceSize()

template<typename Shape , typename UnrolledDescriptorType >

__host__ constexpr __device__ auto Layout< Shape, UnrolledDescriptorType >::GetElementSpaceSize ( ) const

inlineconstexpr

◆ GetLength()

template<typename Shape , typename UnrolledDescriptorType >

template<index_t IDim>

__host__ constexpr __device__ auto Layout< Shape, UnrolledDescriptorType >::GetLength ( ) const

inlineconstexpr

Length getter (product if tuple).

Template Parameters

IDim	Tuple of indexes or index.

Returns: Calculated size.

◆ GetLengths()

template<typename Shape , typename UnrolledDescriptorType >

__host__ constexpr __device__ auto Layout< Shape, UnrolledDescriptorType >::GetLengths ( ) const

inlineconstexpr

Layout size getter (product of shape).

Returns: Calculated size.

◆ GetMergedNestingDescriptor()

template<typename Shape , typename UnrolledDescriptorType >

__host__ constexpr __device__ const MergedNestsDescriptorType& Layout< Shape, UnrolledDescriptorType >::GetMergedNestingDescriptor ( ) const

inlineconstexpr

Get descriptor with all nested dimensions merged. Example, shape: ((2, 2), 2) Descriptor lengths: (4, 2)

Note: The size of merged descriptor is the same as Layout's shape.

Returns: Merged nests descriptor.

◆ GetShape()

template<typename Shape , typename UnrolledDescriptorType >

__host__ constexpr __device__ const Shape& Layout< Shape, UnrolledDescriptorType >::GetShape ( ) const

inlineconstexpr

Shape getter.

Returns: Shape.

◆ GetUnrolledDescriptor()

template<typename Shape , typename UnrolledDescriptorType >

__host__ constexpr __device__ const UnrolledDescriptorType& Layout< Shape, UnrolledDescriptorType >::GetUnrolledDescriptor ( ) const

inlineconstexpr

Get unnested descriptor (with unrolled dims) Example, shape: ((2, 2), 2) Descriptor lengths: (2, 2, 2)

Returns: Flattened descriptor.

◆ operator()() [1/2]

template<typename Shape , typename UnrolledDescriptorType >

template<typename Idxs >

__host__ constexpr __device__ index_t Layout< Shape, UnrolledDescriptorType >::operator() ( ) const

inlineconstexpr

Returns real offset to element in runtime.

Template Parameters

Idxs	Tuple of indexes.

Returns: Calculated offset.

◆ operator()() [2/2]

template<typename Shape , typename UnrolledDescriptorType >

template<typename... Ts>

__host__ __device__ index_t Layout< Shape, UnrolledDescriptorType >::operator() ( const Tuple< Ts... > & Idx ) const

inline

Returns real offset to element in compile time.

Parameters

Idx	Tuple of indexes.

Returns: Calculated offset.

◆ TransformDesc()

template<typename Shape , typename UnrolledDescriptorType >

template<typename... ShapeDims, typename... IdxDims>

__host__ constexpr static __device__ auto Layout< Shape, UnrolledDescriptorType >::TransformDesc	(	const Tuple< ShapeDims... > &	shape,
		const Tuple< IdxDims... > &	idxs,
		const UnrolledDescriptorType &	naive_descriptor
	)

inlinestaticconstexpr

Transform descriptor to align to passed indexes.

Parameters

shape	Layout shape.
idxs	Indexes to align descriptor.
naive_descriptor	Descriptor to merge.

Returns: Aligned descriptor to idx.

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

include/ck/wrapper/layout.hpp

Layout< Shape, UnrolledDescriptorType > Struct Template Reference

Layout&lt; Shape, UnrolledDescriptorType &gt; Struct Template Reference#

Public Types

Public Member Functions

Static Public Member Functions

Detailed Description

template<typename Shape, typename UnrolledDescriptorType> struct Layout< Shape, UnrolledDescriptorType >

Member Typedef Documentation

◆ LayoutShape

◆ LayoutUnrolledDescriptorType

◆ MergedNestsDescriptorType

Constructor & Destructor Documentation

◆ Layout() [1/2]

◆ Layout() [2/2]

Member Function Documentation

◆ Get1DDescriptor()

◆ GetDefaultLengthsTuple()

◆ GetDefaultStartIdxs()

◆ GetElementSpaceSize()

◆ GetLength()

◆ GetLengths()

◆ GetMergedNestingDescriptor()

◆ GetShape()

◆ GetUnrolledDescriptor()

◆ operator()() [1/2]

◆ operator()() [2/2]

◆ TransformDesc()

Layout< Shape, UnrolledDescriptorType > Struct Template Reference#

template<typename Shape, typename UnrolledDescriptorType>
struct Layout< Shape, UnrolledDescriptorType >