Tensor< BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType > Struct Template Reference

Tensor&lt; BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType &gt; Struct Template Reference#

Composable Kernel: Tensor< BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType > Struct Template Reference
Public Types | Public Member Functions | Static Public Attributes | List of all members
Tensor< BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType > Struct Template Reference

Tensor wrapper that performs static and dynamic buffer logic. The tensor is based on a descriptor stored in the Layout. Additionally, tensor can be sliced or shifted using multi-index offset. More...

#include <tensor.hpp>

Public Types

using ElementSpaceSize = decltype(Layout< Shape, UnrolledDescriptorType >{ Shape{}, UnrolledDescriptorType{}}.GetElementSpaceSize())
 
using TensorElementType = std::conditional_t< is_scalar_type< ElementType >::value, ElementType, typename scalar_type< std::remove_const_t< ElementType > >::type >
 

Public Member Functions

__host__ __device__ Tensor ()=delete
 
__host__ constexpr __device__ Tensor (ElementType *pointer, const Layout< Shape, UnrolledDescriptorType > &layout)
 
__host__ constexpr __device__ Tensor (const Layout< Shape, UnrolledDescriptorType > &layout)
 
__host__ constexpr __device__ const Layout< Shape, UnrolledDescriptorType > & GetLayout () const
 
template<typename... Ts, enable_if_t< detail::HasSlice(Tuple< Ts... >{}), bool > = false>
__host__ __device__ auto operator[] (const Tuple< Ts... > &idx)
 Get the new sliced tensor. More...
 
template<typename... Ts, enable_if_t< detail::HasSlice(Tuple< Ts... >{}), bool > = false>
__host__ __device__ auto operator() (const Tuple< Ts... > &idx)
 
template<typename... Idxs, enable_if_t< detail::HasSlice(Tuple< Idxs... >{}), bool > = false>
__host__ __device__ auto operator() (Idxs... idxs)
 
template<typename... Ts, enable_if_t<!detail::HasSlice(Tuple< Ts... >{}), bool > = false>
__host__ __device__ const TensorElementType & operator[] (const Tuple< Ts... > &idx) const
 Getter of the tensor's const value reference. More...
 
template<typename... Ts, enable_if_t<!detail::HasSlice(Tuple< Ts... >{}), bool > = false>
__host__ __device__ const TensorElementType & operator() (const Tuple< Ts... > &idx) const
 
template<typename... Idxs, enable_if_t<!detail::HasSlice(Tuple< Idxs... >{}), bool > = false>
__host__ __device__ const TensorElementType & operator() (Idxs... idxs) const
 
template<typename... Ts, enable_if_t<!detail::HasSlice(Tuple< Ts... >{}), bool > = false>
__host__ __device__ TensorElementType & operator[] (const Tuple< Ts... > &idx)
 Getter of tensor value reference. More...
 
template<typename... Ts, enable_if_t<!detail::HasSlice(Tuple< Ts... >{}), bool > = false>
__host__ __device__ TensorElementType & operator() (const Tuple< Ts... > &idx)
 
template<typename... Idxs, enable_if_t<!detail::HasSlice(Tuple< Idxs... >{}), bool > = false>
__host__ __device__ TensorElementType & operator() (Idxs... idxs)
 
__host__ constexpr __device__ auto GetMergedNestingDescriptor ()
 Get descriptor with all nested dimensions merged. More...
 
__host__ __device__ TensorElementType * GetPointer () const
 Get pointer to the data. More...
 
__host__ constexpr __device__ auto & GetBuffer ()
 
__host__ constexpr __device__ auto & GetBuffer () const
 
__host__ constexpr __device__ auto & GetMultiIdxOffsets () const
 Get multi index offset to the data. More...
 
template<typename MultiIdxOffsets >
__host__ constexpr __device__ void SetMultiIdxOffset (const MultiIdxOffsets multi_idx_offset)
 Apply multi index offset on the tensor. More...
 

Static Public Attributes

static constexpr MemoryTypeEnum TensorBufferAddressSpace = BufferAddressSpace
 
static constexpr bool IsDynamicBuffer
 

Detailed Description

template<MemoryTypeEnum BufferAddressSpace, typename ElementType, typename Shape, typename UnrolledDescriptorType>
struct Tensor< BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType >

Tensor wrapper that performs static and dynamic buffer logic. The tensor is based on a descriptor stored in the Layout. Additionally, tensor can be sliced or shifted using multi-index offset.

Template Parameters
BufferAddressSpaceMemory type (Generic, Global, LDS, VGPR, SGPR).
ElementTypeElement data type.
ShapeTensor shape (layout component).
UnrolledDescriptorTypeFlatten descriptor (layout component).

Member Typedef Documentation

◆ ElementSpaceSize

template<MemoryTypeEnum BufferAddressSpace, typename ElementType , typename Shape , typename UnrolledDescriptorType >
using Tensor< BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType >::ElementSpaceSize = decltype(Layout<Shape, UnrolledDescriptorType>{ Shape{}, UnrolledDescriptorType{}}.GetElementSpaceSize())

◆ TensorElementType

template<MemoryTypeEnum BufferAddressSpace, typename ElementType , typename Shape , typename UnrolledDescriptorType >
using Tensor< BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType >::TensorElementType = std::conditional_t< is_scalar_type<ElementType>::value, ElementType, typename scalar_type<std::remove_const_t<ElementType> >::type>

Constructor & Destructor Documentation

◆ Tensor() [1/3]

template<MemoryTypeEnum BufferAddressSpace, typename ElementType , typename Shape , typename UnrolledDescriptorType >
__host__ __device__ Tensor< BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType >::Tensor ( )
delete

◆ Tensor() [2/3]

template<MemoryTypeEnum BufferAddressSpace, typename ElementType , typename Shape , typename UnrolledDescriptorType >
__host__ constexpr __device__ Tensor< BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType >::Tensor ( ElementType *  pointer,
const Layout< Shape, UnrolledDescriptorType > &  layout 
)
inlineconstexpr

◆ Tensor() [3/3]

template<MemoryTypeEnum BufferAddressSpace, typename ElementType , typename Shape , typename UnrolledDescriptorType >
__host__ constexpr __device__ Tensor< BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType >::Tensor ( const Layout< Shape, UnrolledDescriptorType > &  layout)
inlineconstexpr

Member Function Documentation

◆ GetBuffer() [1/2]

template<MemoryTypeEnum BufferAddressSpace, typename ElementType , typename Shape , typename UnrolledDescriptorType >
__host__ constexpr __device__ auto& Tensor< BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType >::GetBuffer ( )
inlineconstexpr

◆ GetBuffer() [2/2]

template<MemoryTypeEnum BufferAddressSpace, typename ElementType , typename Shape , typename UnrolledDescriptorType >
__host__ constexpr __device__ auto& Tensor< BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType >::GetBuffer ( ) const
inlineconstexpr

◆ GetLayout()

template<MemoryTypeEnum BufferAddressSpace, typename ElementType , typename Shape , typename UnrolledDescriptorType >
__host__ constexpr __device__ const Layout<Shape, UnrolledDescriptorType>& Tensor< BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType >::GetLayout ( ) const
inlineconstexpr

◆ GetMergedNestingDescriptor()

template<MemoryTypeEnum BufferAddressSpace, typename ElementType , typename Shape , typename UnrolledDescriptorType >
__host__ constexpr __device__ auto Tensor< BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType >::GetMergedNestingDescriptor ( )
inlineconstexpr

Get descriptor with all nested dimensions merged.

Returns
Merged nests descriptor.

◆ GetMultiIdxOffsets()

template<MemoryTypeEnum BufferAddressSpace, typename ElementType , typename Shape , typename UnrolledDescriptorType >
__host__ constexpr __device__ auto& Tensor< BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType >::GetMultiIdxOffsets ( ) const
inlineconstexpr

Get multi index offset to the data.

Returns
Multi index offset.

◆ GetPointer()

template<MemoryTypeEnum BufferAddressSpace, typename ElementType , typename Shape , typename UnrolledDescriptorType >
__host__ __device__ TensorElementType* Tensor< BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType >::GetPointer ( ) const
inline

Get pointer to the data.

Returns
Pointer.

◆ operator()() [1/6]

template<MemoryTypeEnum BufferAddressSpace, typename ElementType , typename Shape , typename UnrolledDescriptorType >
template<typename... Ts, enable_if_t< detail::HasSlice(Tuple< Ts... >{}), bool > = false>
__host__ __device__ auto Tensor< BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType >::operator() ( const Tuple< Ts... > &  idx)
inline

◆ operator()() [2/6]

template<MemoryTypeEnum BufferAddressSpace, typename ElementType , typename Shape , typename UnrolledDescriptorType >
template<typename... Ts, enable_if_t<!detail::HasSlice(Tuple< Ts... >{}), bool > = false>
__host__ __device__ TensorElementType& Tensor< BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType >::operator() ( const Tuple< Ts... > &  idx)
inline

◆ operator()() [3/6]

template<MemoryTypeEnum BufferAddressSpace, typename ElementType , typename Shape , typename UnrolledDescriptorType >
template<typename... Ts, enable_if_t<!detail::HasSlice(Tuple< Ts... >{}), bool > = false>
__host__ __device__ const TensorElementType& Tensor< BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType >::operator() ( const Tuple< Ts... > &  idx) const
inline

◆ operator()() [4/6]

template<MemoryTypeEnum BufferAddressSpace, typename ElementType , typename Shape , typename UnrolledDescriptorType >
template<typename... Idxs, enable_if_t< detail::HasSlice(Tuple< Idxs... >{}), bool > = false>
__host__ __device__ auto Tensor< BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType >::operator() ( Idxs...  idxs)
inline

◆ operator()() [5/6]

template<MemoryTypeEnum BufferAddressSpace, typename ElementType , typename Shape , typename UnrolledDescriptorType >
template<typename... Idxs, enable_if_t<!detail::HasSlice(Tuple< Idxs... >{}), bool > = false>
__host__ __device__ TensorElementType& Tensor< BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType >::operator() ( Idxs...  idxs)
inline

◆ operator()() [6/6]

template<MemoryTypeEnum BufferAddressSpace, typename ElementType , typename Shape , typename UnrolledDescriptorType >
template<typename... Idxs, enable_if_t<!detail::HasSlice(Tuple< Idxs... >{}), bool > = false>
__host__ __device__ const TensorElementType& Tensor< BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType >::operator() ( Idxs...  idxs) const
inline

◆ operator[]() [1/3]

template<MemoryTypeEnum BufferAddressSpace, typename ElementType , typename Shape , typename UnrolledDescriptorType >
template<typename... Ts, enable_if_t< detail::HasSlice(Tuple< Ts... >{}), bool > = false>
__host__ __device__ auto Tensor< BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType >::operator[] ( const Tuple< Ts... > &  idx)
inline

Get the new sliced tensor.

Parameters
idxTuple of indices: slice(from,to) or scalar.
Returns
Sliced tensor.

◆ operator[]() [2/3]

template<MemoryTypeEnum BufferAddressSpace, typename ElementType , typename Shape , typename UnrolledDescriptorType >
template<typename... Ts, enable_if_t<!detail::HasSlice(Tuple< Ts... >{}), bool > = false>
__host__ __device__ TensorElementType& Tensor< BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType >::operator[] ( const Tuple< Ts... > &  idx)
inline

Getter of tensor value reference.

Parameters
idxTuple of indices.
Returns
Requested value.

◆ operator[]() [3/3]

template<MemoryTypeEnum BufferAddressSpace, typename ElementType , typename Shape , typename UnrolledDescriptorType >
template<typename... Ts, enable_if_t<!detail::HasSlice(Tuple< Ts... >{}), bool > = false>
__host__ __device__ const TensorElementType& Tensor< BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType >::operator[] ( const Tuple< Ts... > &  idx) const
inline

Getter of the tensor's const value reference.

Parameters
idxTuple of indices.
Returns
Requested value.

◆ SetMultiIdxOffset()

template<MemoryTypeEnum BufferAddressSpace, typename ElementType , typename Shape , typename UnrolledDescriptorType >
template<typename MultiIdxOffsets >
__host__ constexpr __device__ void Tensor< BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType >::SetMultiIdxOffset ( const MultiIdxOffsets  multi_idx_offset)
inlineconstexpr

Apply multi index offset on the tensor.

Parameters
multi_idx_offsetMulti index offset.

Member Data Documentation

◆ IsDynamicBuffer

template<MemoryTypeEnum BufferAddressSpace, typename ElementType , typename Shape , typename UnrolledDescriptorType >
constexpr bool Tensor< BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType >::IsDynamicBuffer
staticconstexpr
Initial value:
= !(BufferAddressSpace == MemoryTypeEnum ::Sgpr ||
BufferAddressSpace == MemoryTypeEnum ::Vgpr)

◆ TensorBufferAddressSpace

template<MemoryTypeEnum BufferAddressSpace, typename ElementType , typename Shape , typename UnrolledDescriptorType >
constexpr MemoryTypeEnum Tensor< BufferAddressSpace, ElementType, Shape, UnrolledDescriptorType >::TensorBufferAddressSpace = BufferAddressSpace
staticconstexpr
The documentation for this struct was generated from the following file:
  • /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/wrapper/tensor.hpp