Cutlass
CUDA Templates for Linear Algebra Subroutines and Solvers
Public Types | Public Member Functions | Static Public Attributes | List of all members
cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ > Class Template Reference

Defines a view into a logical tensor.

#include <tensor_view.h>

Inheritance diagram for cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >:
cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >

Public Types

typedef TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ > Base
 Base tensor reference. More...
 
typedef TensorRef< typename platform::remove_const< Storage_ >::type const, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ > ConstTensorRef
 Tensor reference to of constant value. More...
 
typedef Base TensorRef
 Base tensor reference. More...
 
typedef Base::Storage Storage
 Storage type. More...
 
typedef Base::Index Index
 Index type. More...
 
typedef TensorRef::TensorCoord TensorCoord
 Coordinate in logical tensor space. More...
 
typedef TensorRef::StorageCoord StorageCoord
 Coordinate in storage n-D array. More...
 
typedef TensorRef::StrideVector StrideVector
 
typedef TensorView< typename platform::remove_const< Storage >::type const, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ > ConstTensorView
 TensorView of constant value. More...
 
typedef TensorCoord Coord_t
 Coordinate in logical tensor space. More...
 
typedef Base::LongIndex Offset_t
 Type used to compute the offset of an element to the base of a tensor. More...
 
typedef TensorRef TensorRef_t
 Base class. More...
 
typedef TensorRef::ConstTensorRef ConstTensorRef_t
 TensorRef to const-valued type. More...
 
- Public Types inherited from cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
typedef Storage_ Storage
 Data type of individual access. More...
 
typedef MapFunc_ MapFunc
 Mapping function from logical coordinate to internal n-D array. More...
 
typedef Index_ Index
 Index type. More...
 
typedef LongIndex_ LongIndex
 Typically, strides in memory can be very large. More...
 
typedef Coord< kRankTensorCoord
 Coordinate in logical tensor space. More...
 
typedef Coord< kStorageRankStorageCoord
 Coordinate in storage n-D array. More...
 
typedef Coord< kStorageRank - 1 > StrideVector
 
typedef TensorRef< typename platform::remove_const< Storage >::type const, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ > ConstTensorRef
 Tensor reference to of constant value. More...
 
typedef TensorCoord Coord_t
 Coordinate in logical tensor space. More...
 

Public Member Functions

CUTLASS_HOST_DEVICE TensorView ()
 Default constructor. More...
 
CUTLASS_HOST_DEVICE TensorView (Base const &_ref, TensorCoord const &_size)
 Constructs a TensorView from a TensorRef and size. More...
 
CUTLASS_HOST_DEVICE TensorView (Storage *ptr, StrideVector const &stride, TensorCoord const &size)
 Constructs a TensorView from a pointer, a stride vector, and size. More...
 
CUTLASS_HOST_DEVICE TensorView (Storage *ptr, StorageCoord const &stride, TensorCoord const &size)
 Constructs a TensorView from a pointer, a stride vector, and size. More...
 
CUTLASS_HOST_DEVICE void reset (Base const &_ref=Base(), TensorCoord const &_size=TensorCoord())
 Updates the reference and size of a Tensor_view object. More...
 
CUTLASS_HOST_DEVICE TensorCoord const & size () const
 Accesses the size. More...
 
CUTLASS_HOST_DEVICE Index size (int dim) const
 Accesses the size. More...
 
CUTLASS_HOST_DEVICE TensorViewoperator= (TensorView const &_tensor)
 Assigns the Tensor_view. More...
 
CUTLASS_HOST_DEVICE bool contains (TensorCoord const &coord) const
 Determines whether a location is within a tensor. More...
 
CUTLASS_HOST_DEVICE TensorRef ref () const
 Returns a TensorRef pointing to the first element of the tensor. More...
 
CUTLASS_HOST_DEVICE ConstTensorRef const_ref () const
 Returns a TensorRef pointing to the first element of the tensor. More...
 
CUTLASS_HOST_DEVICE TensorView subview (TensorCoord const &location, TensorCoord size) const
 Returns a Tensor_view given location and size quantities. More...
 
CUTLASS_HOST_DEVICE size_t capacity () const
 Returns the number of scalar elements needed to store tensor. More...
 
CUTLASS_HOST_DEVICE TensorView operator+ (TensorCoord const &b) const
 Returns a TensorView offset by a given amount. More...
 
CUTLASS_HOST_DEVICE TensorViewoperator+= (TensorCoord const &b)
 Returns a TensorRef offset by a given amount. More...
 
CUTLASS_HOST_DEVICE TensorView operator- (TensorCoord const &b) const
 Returns a TensorRef offset by a given amount. More...
 
CUTLASS_HOST_DEVICE TensorViewoperator-= (TensorCoord const &b)
 Returns a TensorRef offset by a given amount. More...
 
- Public Member Functions inherited from cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
CUTLASS_HOST_DEVICE TensorRef (Storage *ptr=nullptr)
 Helper for 1-D memory. All higher ranks are projected onto the fastest changing rank. More...
 
CUTLASS_HOST_DEVICE TensorRef (Storage *ptr, Index ldm)
 Helper to construct from a pointer and single stride element for 2-D pitch linear memory. More...
 
CUTLASS_HOST_DEVICE TensorRef (Storage *ptr, StrideVector const &stride)
 Constructs from a single pointer and stride vector. More...
 
CUTLASS_HOST_DEVICE TensorRef (Storage *ptr, StorageCoord const &stride)
 
CUTLASS_HOST_DEVICE TensorRef (TensorRef< typename platform::remove_const< Storage >::type, kRank, MapFunc, kStorageRank, Index, LongIndex > const &ref)
 Enables conversion from TensorRef of non-const type. More...
 
CUTLASS_HOST_DEVICE ConstTensorRef const_ref () const
 Returns a reference to constant-valued tensor. More...
 
CUTLASS_HOST_DEVICE void reset (Storage *ptr=nullptr)
 Updates only the pointer. More...
 
CUTLASS_HOST_DEVICE void reset (Storage *ptr, StorageCoord const &stride)
 Updates the pointer, stride, and location within a TensorRef. More...
 
CUTLASS_HOST_DEVICE bool good () const
 Returns true if the TensorRef may be safely accessed. More...
 
CUTLASS_HOST_DEVICE Storagedata () const
 Returns the pointer to referenced data. More...
 
CUTLASS_HOST_DEVICE StorageCoord stride () const
 Returns the stride of the tensor. More...
 
CUTLASS_HOST_DEVICE Index stride (int dim) const
 Returns the stride of the tensor in the given dimension. More...
 
CUTLASS_HOST_DEVICE Index leading_dim (int idx=0) const
 Returns the maximum stride element as the 'leading dimension'. More...
 
CUTLASS_HOST_DEVICE StorageCoord map (TensorCoord const &coord) const
 Maps a logical coordinate to an n-D array in memory. More...
 
CUTLASS_HOST_DEVICE LongIndex offset (TensorCoord const &coord) const
 Computes the offset of an index from the origin of the tensor. More...
 
CUTLASS_HOST_DEVICE Storageat (TensorCoord const &coord) const
 Returns a reference to the element at a given Coord. More...
 
CUTLASS_HOST_DEVICE Storageat (LongIndex idx) const
 Returns a reference to the element at a given linear index. More...
 
CUTLASS_HOST_DEVICE Storageoperator[] (TensorCoord const &coord) const
 Returns a reference to the element at a given Coord. More...
 
CUTLASS_HOST_DEVICE Storageoperator[] (LongIndex idx) const
 Returns a reference to the element at a given linear index. More...
 
CUTLASS_HOST_DEVICE TensorRefadd_pointer_offset (LongIndex delta)
 Adds an offset to each pointer. More...
 
CUTLASS_HOST_DEVICE TensorRef operator+ (TensorCoord const &b) const
 Returns a TensorRef offset by a given amount. More...
 
CUTLASS_HOST_DEVICE TensorRefoperator+= (TensorCoord const &b)
 Returns a TensorRef offset by a given amount. More...
 
CUTLASS_HOST_DEVICE TensorRef operator- (TensorCoord const &b) const
 Returns a TensorRef offset by a given amount. More...
 
CUTLASS_HOST_DEVICE TensorRefoperator-= (TensorCoord const &b)
 Returns a TensorRef offset by a given amount. More...
 

Static Public Attributes

static int const Rank = Base::kRank
 Logical rank of tensor index space. More...
 
- Static Public Attributes inherited from cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
static int const kRank = Rank_
 Logical rank of tensor index space. More...
 
static int const kStorageRank = StorageRank_
 Rank of internal storage. More...
 
static int const Rank = kRank
 Logical rank of tensor index space. More...
 

Member Typedef Documentation

◆ Base

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
typedef TensorRef<Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_> cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::Base

◆ ConstTensorRef

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
typedef TensorRef< typename platform::remove_const<Storage_>::type const, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_> cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstTensorRef

◆ ConstTensorRef_t

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
typedef TensorRef::ConstTensorRef cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstTensorRef_t

◆ ConstTensorView

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
typedef TensorView< typename platform::remove_const<Storage>::type const, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_> cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstTensorView

◆ Coord_t

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
typedef TensorCoord cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::Coord_t

◆ Index

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
typedef Base::Index cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::Index

◆ Offset_t

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
typedef Base::LongIndex cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::Offset_t

◆ Storage

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
typedef Base::Storage cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::Storage

◆ StorageCoord

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
typedef TensorRef::StorageCoord cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::StorageCoord

◆ StrideVector

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
typedef TensorRef::StrideVector cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::StrideVector

Stride vector in storage coordinate space Least significant stride is = 1 and not stored

◆ TensorCoord

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
typedef TensorRef::TensorCoord cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::TensorCoord

◆ TensorRef

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
typedef Base cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::TensorRef

◆ TensorRef_t

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
typedef TensorRef cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::TensorRef_t

Constructor & Destructor Documentation

◆ TensorView() [1/4]

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
CUTLASS_HOST_DEVICE cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::TensorView ( )
inline

◆ TensorView() [2/4]

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
CUTLASS_HOST_DEVICE cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::TensorView ( Base const &  _ref,
TensorCoord const &  _size 
)
inline

◆ TensorView() [3/4]

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
CUTLASS_HOST_DEVICE cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::TensorView ( Storage ptr,
StrideVector const &  stride,
TensorCoord const &  size 
)
inline

◆ TensorView() [4/4]

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
CUTLASS_HOST_DEVICE cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::TensorView ( Storage ptr,
StorageCoord const &  stride,
TensorCoord const &  size 
)
inline

Member Function Documentation

◆ capacity()

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
CUTLASS_HOST_DEVICE size_t cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::capacity ( ) const
inline

◆ const_ref()

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
CUTLASS_HOST_DEVICE ConstTensorRef cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::const_ref ( ) const
inline

◆ contains()

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
CUTLASS_HOST_DEVICE bool cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::contains ( TensorCoord const &  coord) const
inline

◆ operator+()

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
CUTLASS_HOST_DEVICE TensorView cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::operator+ ( TensorCoord const &  b) const
inline

◆ operator+=()

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
CUTLASS_HOST_DEVICE TensorView& cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::operator+= ( TensorCoord const &  b)
inline

◆ operator-()

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
CUTLASS_HOST_DEVICE TensorView cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::operator- ( TensorCoord const &  b) const
inline

◆ operator-=()

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
CUTLASS_HOST_DEVICE TensorView& cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::operator-= ( TensorCoord const &  b)
inline

◆ operator=()

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
CUTLASS_HOST_DEVICE TensorView& cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::operator= ( TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ > const &  _tensor)
inline

◆ ref()

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
CUTLASS_HOST_DEVICE TensorRef cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ref ( ) const
inline

◆ reset()

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
CUTLASS_HOST_DEVICE void cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::reset ( Base const &  _ref = Base(),
TensorCoord const &  _size = TensorCoord() 
)
inline

◆ size() [1/2]

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
CUTLASS_HOST_DEVICE TensorCoord const& cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::size ( ) const
inline

◆ size() [2/2]

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
CUTLASS_HOST_DEVICE Index cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::size ( int  dim) const
inline

◆ subview()

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
CUTLASS_HOST_DEVICE TensorView cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::subview ( TensorCoord const &  location,
TensorCoord  size 
) const
inline

Member Data Documentation

◆ Rank

template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
int const cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::Rank = Base::kRank
static

The documentation for this class was generated from the following file: