Tensor< T > Struct Template Reference

Tensor&lt; T &gt; Struct Template Reference#

Composable Kernel: Tensor< T > 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 <host_tensor.hpp>

Public Types

using Descriptor = HostTensorDescriptor
 
using Data = std::vector< T >
 
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

template<typename X >
 Tensor (std::initializer_list< X > lens)
 
template<typename X , typename Y >
 Tensor (std::initializer_list< X > lens, std::initializer_list< Y > strides)
 
template<typename Lengths >
 Tensor (const Lengths &lens)
 
template<typename Lengths , typename Strides >
 Tensor (const Lengths &lens, const Strides &strides)
 
 Tensor (const Descriptor &desc)
 
template<typename OutT >
Tensor< OutT > CopyAsType () const
 
 Tensor ()=delete
 
 Tensor (const Tensor &)=default
 
 Tensor (Tensor &&)=default
 
 ~Tensor ()=default
 
Tensoroperator= (const Tensor &)=default
 
Tensoroperator= (Tensor &&)=default
 
template<typename FromT >
 Tensor (const Tensor< FromT > &other)
 
void savetxt (std::string file_name, std::string dtype="float")
 
decltype(auto) GetLengths () const
 
decltype(auto) GetStrides () const
 
std::size_t GetNumOfDimension () const
 
std::size_t GetElementSize () const
 
std::size_t GetElementSpaceSize () const
 
std::size_t GetElementSpaceSizeInBytes () const
 
void SetZero ()
 
template<typename F >
void ForEach_impl (F &&f, std::vector< size_t > &idx, size_t rank)
 
template<typename F >
void ForEach (F &&f)
 
template<typename F >
void ForEach_impl (const F &&f, std::vector< size_t > &idx, size_t rank) const
 
template<typename F >
void ForEach (const F &&f) const
 
template<typename G >
void GenerateTensorValue (G g, std::size_t num_thread=1)
 
template<typename Distribution = std::uniform_real_distribution<float>, typename Mapping = ck::identity, typename Generator = std::minstd_rand>
void GenerateTensorDistr (Distribution dis={0.f, 1.f}, Mapping fn={}, const Generator g=Generator(0), std::size_t num_thread=-1)
 
template<typename... Is>
std::size_t GetOffsetFromMultiIndex (Is... is) const
 
template<typename... Is>
T & operator() (Is... is)
 
template<typename... Is>
const T & operator() (Is... is) const
 
T & operator() (const std::vector< std::size_t > &idx)
 
const T & operator() (const std::vector< std::size_t > &idx) const
 
Data::iterator begin ()
 
Data::iterator end ()
 
Data::pointer data ()
 
Data::const_iterator begin () const
 
Data::const_iterator end () const
 
Data::const_pointer data () const
 
Data::size_type size () const
 
template<typename U = T>
auto AsSpan () const
 
template<typename U = T>
auto AsSpan ()
 
__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 TensorElementTypeoperator[] (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 TensorElementTypeoperator() (const Tuple< Ts... > &idx) const
 
template<typename... Idxs, enable_if_t<!detail::HasSlice(Tuple< Idxs... >{}), bool > = false>
__host__ __device__ const TensorElementTypeoperator() (Idxs... idxs) const
 
template<typename... Ts, enable_if_t<!detail::HasSlice(Tuple< Ts... >{}), bool > = false>
__host__ __device__ TensorElementTypeoperator[] (const Tuple< Ts... > &idx)
 Getter of tensor value reference. More...
 
template<typename... Ts, enable_if_t<!detail::HasSlice(Tuple< Ts... >{}), bool > = false>
__host__ __device__ TensorElementTypeoperator() (const Tuple< Ts... > &idx)
 
template<typename... Idxs, enable_if_t<!detail::HasSlice(Tuple< Idxs... >{}), bool > = false>
__host__ __device__ TensorElementTypeoperator() (Idxs... idxs)
 
__host__ constexpr __device__ auto GetMergedNestingDescriptor ()
 Get descriptor with all nested dimensions merged. More...
 
__host__ __device__ TensorElementTypeGetPointer () 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...
 

Public Attributes

Descriptor mDesc
 
Data mData
 

Static Public Attributes

static constexpr MemoryTypeEnum TensorBufferAddressSpace = BufferAddressSpace
 
static constexpr bool IsDynamicBuffer
 

Detailed Description

template<typename T>
struct Tensor< T >

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

◆ Data

template<typename T >
using Tensor< T >::Data = std::vector<T>

◆ Descriptor

template<typename T >
using Tensor< T >::Descriptor = HostTensorDescriptor

◆ ElementSpaceSize

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

◆ TensorElementType

template<typename T >
using Tensor< T >::TensorElementType = std::conditional_t< is_scalar_type<ElementType>::value, ElementType, typename scalar_type<std::remove_const_t<ElementType> >::type>

Constructor & Destructor Documentation

◆ Tensor() [1/12]

template<typename T >
template<typename X >
Tensor< T >::Tensor ( std::initializer_list< X >  lens)
inline

◆ Tensor() [2/12]

template<typename T >
template<typename X , typename Y >
Tensor< T >::Tensor ( std::initializer_list< X >  lens,
std::initializer_list< Y >  strides 
)
inline

◆ Tensor() [3/12]

template<typename T >
template<typename Lengths >
Tensor< T >::Tensor ( const Lengths &  lens)
inline

◆ Tensor() [4/12]

template<typename T >
template<typename Lengths , typename Strides >
Tensor< T >::Tensor ( const Lengths &  lens,
const Strides &  strides 
)
inline

◆ Tensor() [5/12]

template<typename T >
Tensor< T >::Tensor ( const Descriptor desc)
inline

◆ Tensor() [6/12]

template<typename T >
Tensor< T >::Tensor ( )
delete

◆ Tensor() [7/12]

template<typename T >
Tensor< T >::Tensor ( const Tensor< T > &  )
default

◆ Tensor() [8/12]

template<typename T >
Tensor< T >::Tensor ( Tensor< T > &&  )
default

◆ ~Tensor()

template<typename T >
Tensor< T >::~Tensor ( )
default

◆ Tensor() [9/12]

template<typename T >
template<typename FromT >
Tensor< T >::Tensor ( const Tensor< FromT > &  other)
inlineexplicit

◆ Tensor() [10/12]

template<typename T >
__host__ __device__ Tensor< T >::Tensor ( )
delete

◆ Tensor() [11/12]

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

◆ Tensor() [12/12]

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

Member Function Documentation

◆ AsSpan() [1/2]

template<typename T >
template<typename U = T>
auto Tensor< T >::AsSpan ( )
inline

◆ AsSpan() [2/2]

template<typename T >
template<typename U = T>
auto Tensor< T >::AsSpan ( ) const
inline

◆ begin() [1/2]

template<typename T >
Data::iterator Tensor< T >::begin ( )
inline

◆ begin() [2/2]

template<typename T >
Data::const_iterator Tensor< T >::begin ( ) const
inline

◆ CopyAsType()

template<typename T >
template<typename OutT >
Tensor<OutT> Tensor< T >::CopyAsType ( ) const
inline

◆ data() [1/2]

template<typename T >
Data::pointer Tensor< T >::data ( )
inline

◆ data() [2/2]

template<typename T >
Data::const_pointer Tensor< T >::data ( ) const
inline

◆ end() [1/2]

template<typename T >
Data::iterator Tensor< T >::end ( )
inline

◆ end() [2/2]

template<typename T >
Data::const_iterator Tensor< T >::end ( ) const
inline

◆ ForEach() [1/2]

template<typename T >
template<typename F >
void Tensor< T >::ForEach ( const F &&  f) const
inline

◆ ForEach() [2/2]

template<typename T >
template<typename F >
void Tensor< T >::ForEach ( F &&  f)
inline

◆ ForEach_impl() [1/2]

template<typename T >
template<typename F >
void Tensor< T >::ForEach_impl ( const F &&  f,
std::vector< size_t > &  idx,
size_t  rank 
) const
inline

◆ ForEach_impl() [2/2]

template<typename T >
template<typename F >
void Tensor< T >::ForEach_impl ( F &&  f,
std::vector< size_t > &  idx,
size_t  rank 
)
inline

◆ GenerateTensorDistr()

template<typename T >
template<typename Distribution = std::uniform_real_distribution<float>, typename Mapping = ck::identity, typename Generator = std::minstd_rand>
void Tensor< T >::GenerateTensorDistr ( Distribution  dis = {0.f, 1.f},
Mapping  fn = {},
const Generator  g = Generator(0),
std::size_t  num_thread = -1 
)
inline

◆ GenerateTensorValue()

template<typename T >
template<typename G >
void Tensor< T >::GenerateTensorValue ( g,
std::size_t  num_thread = 1 
)
inline

◆ GetBuffer() [1/2]

template<typename T >
__host__ constexpr __device__ auto& Tensor< T >::GetBuffer ( )
inlineconstexpr

◆ GetBuffer() [2/2]

template<typename T >
__host__ constexpr __device__ auto& Tensor< T >::GetBuffer ( ) const
inlineconstexpr

◆ GetElementSize()

template<typename T >
std::size_t Tensor< T >::GetElementSize ( ) const
inline

◆ GetElementSpaceSize()

template<typename T >
std::size_t Tensor< T >::GetElementSpaceSize ( ) const
inline

◆ GetElementSpaceSizeInBytes()

template<typename T >
std::size_t Tensor< T >::GetElementSpaceSizeInBytes ( ) const
inline

◆ GetLayout()

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

◆ GetLengths()

template<typename T >
decltype(auto) Tensor< T >::GetLengths ( ) const
inline

◆ GetMergedNestingDescriptor()

template<typename T >
__host__ constexpr __device__ auto Tensor< T >::GetMergedNestingDescriptor ( )
inlineconstexpr

Get descriptor with all nested dimensions merged.

Returns
Merged nests descriptor.

◆ GetMultiIdxOffsets()

template<typename T >
__host__ constexpr __device__ auto& Tensor< T >::GetMultiIdxOffsets ( ) const
inlineconstexpr

Get multi index offset to the data.

Returns
Multi index offset.

◆ GetNumOfDimension()

template<typename T >
std::size_t Tensor< T >::GetNumOfDimension ( ) const
inline

◆ GetOffsetFromMultiIndex()

template<typename T >
template<typename... Is>
std::size_t Tensor< T >::GetOffsetFromMultiIndex ( Is...  is) const
inline

◆ GetPointer()

template<typename T >
__host__ __device__ TensorElementType* Tensor< T >::GetPointer ( ) const
inline

Get pointer to the data.

Returns
Pointer.

◆ GetStrides()

template<typename T >
decltype(auto) Tensor< T >::GetStrides ( ) const
inline

◆ operator()() [1/10]

template<typename T >
T& Tensor< T >::operator() ( const std::vector< std::size_t > &  idx)
inline

◆ operator()() [2/10]

template<typename T >
const T& Tensor< T >::operator() ( const std::vector< std::size_t > &  idx) const
inline

◆ operator()() [3/10]

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

◆ operator()() [4/10]

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

◆ operator()() [5/10]

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

◆ operator()() [6/10]

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

◆ operator()() [7/10]

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

◆ operator()() [8/10]

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

◆ operator()() [9/10]

template<typename T >
template<typename... Is>
T& Tensor< T >::operator() ( Is...  is)
inline

◆ operator()() [10/10]

template<typename T >
template<typename... Is>
const T& Tensor< T >::operator() ( Is...  is) const
inline

◆ operator=() [1/2]

template<typename T >
Tensor& Tensor< T >::operator= ( const Tensor< T > &  )
default

◆ operator=() [2/2]

template<typename T >
Tensor& Tensor< T >::operator= ( Tensor< T > &&  )
default

◆ operator[]() [1/3]

template<typename T >
template<typename... Ts, enable_if_t< detail::HasSlice(Tuple< Ts... >{}), bool > = false>
__host__ __device__ auto Tensor< T >::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<typename T >
template<typename... Ts, enable_if_t<!detail::HasSlice(Tuple< Ts... >{}), bool > = false>
__host__ __device__ TensorElementType& Tensor< T >::operator[] ( const Tuple< Ts... > &  idx)
inline

Getter of tensor value reference.

Parameters
idxTuple of indices.
Returns
Requested value.

◆ operator[]() [3/3]

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

Getter of the tensor's const value reference.

Parameters
idxTuple of indices.
Returns
Requested value.

◆ savetxt()

template<typename T >
void Tensor< T >::savetxt ( std::string  file_name,
std::string  dtype = "float" 
)
inline

◆ SetMultiIdxOffset()

template<typename T >
template<typename MultiIdxOffsets >
__host__ constexpr __device__ void Tensor< T >::SetMultiIdxOffset ( const MultiIdxOffsets  multi_idx_offset)
inlineconstexpr

Apply multi index offset on the tensor.

Parameters
multi_idx_offsetMulti index offset.

◆ SetZero()

template<typename T >
void Tensor< T >::SetZero ( )
inline

◆ size()

template<typename T >
Data::size_type Tensor< T >::size ( ) const
inline

Member Data Documentation

◆ IsDynamicBuffer

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

◆ mData

template<typename T >
Data Tensor< T >::mData

◆ mDesc

template<typename T >
Descriptor Tensor< T >::mDesc

◆ TensorBufferAddressSpace

template<typename T >
constexpr MemoryTypeEnum Tensor< T >::TensorBufferAddressSpace = BufferAddressSpace
staticconstexpr

The documentation for this struct was generated from the following files:
  • /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/library/utility/host_tensor.hpp
  • /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/wrapper/tensor.hpp