Tensor< T > Struct Template Reference

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

Composable Kernel: Tensor< T > Struct Template Reference
Public Types | Public Member Functions | Public Attributes | Static Public Attributes | List of all members
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
using TensorElementType

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)
template<typename X, typename... Rest, std::enable_if_t<(sizeof...(Rest) > 0), int > = 0>
 Tensor (std::initializer_list< X > lens, Rest &&... rest)
template<typename X, typename Y, typename... Rest, std::enable_if_t<(sizeof...(Rest) > 0), int > = 0>
 Tensor (std::initializer_list< X > lens, std::initializer_list< Y > strides, Rest &&... rest)
template<typename Lengths, typename... Rest, std::enable_if_t<(sizeof...(Rest) > 0), int > = 0>
 Tensor (const Lengths &lens, Rest &&... rest)
template<typename Lengths, typename Strides, typename... Rest, std::enable_if_t<(sizeof...(Rest) > 0), int > = 0>
 Tensor (const Lengths &lens, const Strides &strides, Rest &&... rest)
 Tensor (const Descriptor &desc)
template<typename OutT>
Tensor< OutT > CopyAsType () const
 Tensor ()=delete
 Tensor (const Tensor &)=default
 Tensor (Tensor &&)=default
 ~Tensor ()=default
Tensor & operator= (const Tensor &)=default
Tensor & operator= (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__ __device__ constexpr Tensor (ElementType *pointer, const Layout< Shape, UnrolledDescriptorType > &layout)
__host__ __device__ constexpr Tensor (const Layout< Shape, UnrolledDescriptorType > &layout)
__host__ __device__ constexpr 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.
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.
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.
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__ __device__ constexpr auto GetMergedNestingDescriptor ()
 Get descriptor with all nested dimensions merged.
__host__ __device__ TensorElementType * GetPointer () const
 Get pointer to the data.
__host__ __device__ constexpr auto & GetBuffer ()
__host__ __device__ constexpr auto & GetBuffer () const
__host__ __device__ constexpr auto & GetMultiIdxOffsets () const
 Get multi index offset to the data.
template<typename MultiIdxOffsets>
__host__ __device__ constexpr void SetMultiIdxOffset (const MultiIdxOffsets multi_idx_offset)
 Apply multi index offset on the tensor.

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
Initial value:
decltype(Layout<Shape, UnrolledDescriptorType>{
Shape{}, UnrolledDescriptorType{}}.GetElementSpaceSize())
Layout
Layout wrapper that performs the tensor descriptor logic.
Definition layout.hpp:24
Layout::GetElementSpaceSize
__host__ __device__ constexpr auto GetElementSpaceSize() const
Definition layout.hpp:297

◆ TensorElementType

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

Constructor & Destructor Documentation

◆ Tensor() [1/16]

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

◆ Tensor() [2/16]

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

◆ Tensor() [3/16]

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

◆ Tensor() [4/16]

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

◆ Tensor() [5/16]

template<typename T>
template<typename X, typename... Rest, std::enable_if_t<(sizeof...(Rest) > 0), int > = 0>
Tensor< T >::Tensor ( std::initializer_list< X > lens,
Rest &&... rest )
inline

◆ Tensor() [6/16]

template<typename T>
template<typename X, typename Y, typename... Rest, std::enable_if_t<(sizeof...(Rest) > 0), int > = 0>
Tensor< T >::Tensor ( std::initializer_list< X > lens,
std::initializer_list< Y > strides,
Rest &&... rest )
inline

◆ Tensor() [7/16]

template<typename T>
template<typename Lengths, typename... Rest, std::enable_if_t<(sizeof...(Rest) > 0), int > = 0>
Tensor< T >::Tensor ( const Lengths & lens,
Rest &&... rest )
inline

◆ Tensor() [8/16]

template<typename T>
template<typename Lengths, typename Strides, typename... Rest, std::enable_if_t<(sizeof...(Rest) > 0), int > = 0>
Tensor< T >::Tensor ( const Lengths & lens,
const Strides & strides,
Rest &&... rest )
inline

◆ Tensor() [9/16]

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

◆ Tensor() [10/16]

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

◆ Tensor() [11/16]

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

◆ Tensor() [12/16]

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

◆ ~Tensor()

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

◆ Tensor() [13/16]

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

◆ Tensor() [14/16]

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

◆ Tensor() [15/16]

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

◆ Tensor() [16/16]

template<typename T>
__host__ __device__ constexpr 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 g,
std::size_t num_thread = 1 )
inline

◆ GetBuffer() [1/2]

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

◆ GetBuffer() [2/2]

template<typename T>
__host__ __device__ constexpr 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__ __device__ constexpr 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__ __device__ constexpr auto Tensor< T >::GetMergedNestingDescriptor ( )
inlineconstexpr

Get descriptor with all nested dimensions merged.

Returns
Merged nests descriptor.

◆ GetMultiIdxOffsets()

template<typename T>
__host__ __device__ constexpr 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__ TensorElementType & 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__ auto 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__ TensorElementType & 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__ auto 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__ TensorElementType & Tensor< T >::operator[] ( const Tuple< Ts... > & idx)
inline

Getter of tensor value reference.

Parameters
idxTuple of indices.
Returns
Requested value.

◆ operator[]() [2/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[]() [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__ __device__ constexpr 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>
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>
MemoryTypeEnum Tensor< T >::TensorBufferAddressSpace = BufferAddressSpace
staticconstexpr
The documentation for this struct was generated from the following files: