Tensor< T > Struct Template Reference

Tensor< T > 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)

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__ 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...

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

BufferAddressSpace	Memory type (Generic, Global, LDS, VGPR, SGPR).
ElementType	Element data type.
Shape	Tensor shape (layout component).
UnrolledDescriptorType	Flatten 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/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__ constexpr __device__ Tensor< T >::Tensor	(	ElementType *	pointer,
		const Layout< Shape, UnrolledDescriptorType > &	layout
	)

inlineconstexpr

◆ Tensor() [16/16]

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	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

idx	Tuple 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

idx	Tuple 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

idx	Tuple 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_offset Multi 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

Tensor< T > Struct Template Reference

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

Public Types

Public Member Functions

Public Attributes

Static Public Attributes

Detailed Description

template<typename T> struct Tensor< T >

Member Typedef Documentation

◆ Data

◆ Descriptor

◆ ElementSpaceSize

◆ TensorElementType

Constructor & Destructor Documentation

◆ Tensor() [1/16]

◆ Tensor() [2/16]

◆ Tensor() [3/16]

◆ Tensor() [4/16]

◆ Tensor() [5/16]

◆ Tensor() [6/16]

◆ Tensor() [7/16]

◆ Tensor() [8/16]

◆ Tensor() [9/16]

◆ Tensor() [10/16]

◆ Tensor() [11/16]

◆ Tensor() [12/16]

◆ ~Tensor()

◆ Tensor() [13/16]

◆ Tensor() [14/16]

◆ Tensor() [15/16]

◆ Tensor() [16/16]

Member Function Documentation

◆ AsSpan() [1/2]

◆ AsSpan() [2/2]

◆ begin() [1/2]

◆ begin() [2/2]

◆ CopyAsType()

◆ data() [1/2]

◆ data() [2/2]

◆ end() [1/2]

◆ end() [2/2]

◆ ForEach() [1/2]

◆ ForEach() [2/2]

◆ ForEach_impl() [1/2]

◆ ForEach_impl() [2/2]

◆ GenerateTensorDistr()

◆ GenerateTensorValue()

◆ GetBuffer() [1/2]

◆ GetBuffer() [2/2]

◆ GetElementSize()

◆ GetElementSpaceSize()

◆ GetElementSpaceSizeInBytes()

◆ GetLayout()

◆ GetLengths()

◆ GetMergedNestingDescriptor()

◆ GetMultiIdxOffsets()

◆ GetNumOfDimension()

◆ GetOffsetFromMultiIndex()

◆ GetPointer()

◆ GetStrides()

◆ operator()() [1/10]

◆ operator()() [2/10]

◆ operator()() [3/10]

◆ operator()() [4/10]

◆ operator()() [5/10]

◆ operator()() [6/10]

◆ operator()() [7/10]

◆ operator()() [8/10]

◆ operator()() [9/10]

◆ operator()() [10/10]

◆ operator=() [1/2]

◆ operator=() [2/2]

◆ operator[]() [1/3]

◆ operator[]() [2/3]

◆ operator[]() [3/3]

◆ savetxt()

◆ SetMultiIdxOffset()

◆ SetZero()

◆ size()

Member Data Documentation

Tensor< T > Struct Template Reference#

template<typename T>
struct Tensor< T >