Layout< Shape, UnrolledDescriptorType > Struct Template Reference

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

Composable Kernel: Layout< Shape, UnrolledDescriptorType > Struct Template Reference
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 MergedNestsDescriptorTypeGetMergedNestingDescriptor () 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
ShapeTuple 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.
UnrolledDescriptorTypeTensor 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
shapeShape for layout.
unnested_descriptorDescriptor

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
IDimTuple 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
IdxsTuple 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
IdxTuple 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
shapeLayout shape.
idxsIndexes to align descriptor.
naive_descriptorDescriptor to merge.
Returns
Aligned descriptor to idx.

The documentation for this struct was generated from the following file:
  • /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-6.4.3/include/ck/wrapper/layout.hpp