Cutlass CUDA Templates for Linear Algebra Subroutines and Solvers

Public Types | Public Member Functions | Static Public Attributes | List of all members
cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ > Class Template Reference

`#include <tensor_ref.h>`

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

## Public Types
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< kRank > | TensorCoord
| Coordinate in logical tensor space. More...
typedef Coord< kStorageRank > | StorageCoord
| 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 | 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 Storage * | data () 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 Storage & | at (TensorCoord const &coord) const
| Returns a reference to the element at a given Coord. More...
CUTLASS_HOST_DEVICE Storage & | at (LongIndex idx) const
| Returns a reference to the element at a given linear index. More...
CUTLASS_HOST_DEVICE Storage & | operator[] (TensorCoord const &coord) const
| Returns a reference to the element at a given Coord. More...
CUTLASS_HOST_DEVICE Storage & | operator[] (LongIndex idx) const
| Returns a reference to the element at a given linear index. More...
CUTLASS_HOST_DEVICE TensorRef & | add_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 TensorRef & | operator+= (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 TensorRef & | operator-= (TensorCoord const &b)
| Returns a TensorRef offset by a given amount. More...

## Static Public Attributes
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

## ◆ ConstTensorRef

template<typename Storage_, int Rank_, 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::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstTensorRef

## ◆ Coord_t

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

typedef TensorCoord cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::Coord_t

Require at least rank=1. Mathematically, a rank=0 tensor would be considered to be a scalar, but degenerate cases such as these are difficult to accommodate without extensive C++ metaprogramming or support for zero-length arrays.

## ◆ Index

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

typedef Index_ cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::Index

## ◆ LongIndex

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

typedef LongIndex_ cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::LongIndex

## ◆ MapFunc

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

typedef MapFunc_ cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::MapFunc

## ◆ Storage

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

typedef Storage_ cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::Storage

## ◆ StorageCoord

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

typedef Coord<kStorageRank> cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::StorageCoord

## ◆ StrideVector

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

typedef Coord<kStorageRank - 1> cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::StrideVector

Stride vector in storage coordinage space - assumes least significant stride is 1 and does not store it.

## ◆ TensorCoord

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

typedef Coord<kRank> cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::TensorCoord

## Constructor & Destructor Documentation

## ◆ TensorRef() [1/5]

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

CUTLASS_HOST_DEVICE cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::TensorRef | ( | Storage * | ptr = `nullptr` | ) |
| inline

## ◆ TensorRef() [2/5]

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

CUTLASS_HOST_DEVICE cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::TensorRef | ( | Storage * | ptr,
|  | Index | ldm
| ) |  |
| inline

## ◆ TensorRef() [3/5]

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

CUTLASS_HOST_DEVICE cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::TensorRef | ( | Storage * | ptr,
|  | StrideVector const & | stride
| ) |  |
| inline

## ◆ TensorRef() [4/5]

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

CUTLASS_HOST_DEVICE cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::TensorRef | ( | Storage * | ptr,
|  | StorageCoord const & | stride
| ) |  |
| inline

Constructs from a pointer and a stride vector of size kRank. If fastest changing stride is not 1, construction fails and subsequent calls to good() will return false.

## ◆ TensorRef() [5/5]

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

CUTLASS_HOST_DEVICE cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::TensorRef | ( | TensorRef< typename platform::remove_const< Storage >::type, kRank, MapFunc, kStorageRank, Index, LongIndex > const & | ref | ) |
| inline

## Member Function Documentation

## ◆ add_pointer_offset()

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

CUTLASS_HOST_DEVICE TensorRef& cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::add_pointer_offset | ( | LongIndex | delta | ) |
| inline

## ◆ at() [1/2]

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

CUTLASS_HOST_DEVICE Storage& cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::at | ( | TensorCoord const & | coord | ) | const
| inline

## ◆ at() [2/2]

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

CUTLASS_HOST_DEVICE Storage& cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::at | ( | LongIndex | idx | ) | const
| inline

## ◆ const_ref()

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

CUTLASS_HOST_DEVICE ConstTensorRef cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::const_ref | ( |  | ) | const
| inline

## ◆ data()

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

CUTLASS_HOST_DEVICE Storage* cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::data | ( |  | ) | const
| inline

## ◆ good()

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

CUTLASS_HOST_DEVICE bool cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::good | ( |  | ) | const
| inline

## ◆ leading_dim()

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

CUTLASS_HOST_DEVICE Index cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::leading_dim | ( | int | idx = `0` | ) | const
| inline

## ◆ map()

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

CUTLASS_HOST_DEVICE StorageCoord cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::map | ( | TensorCoord const & | coord | ) | const
| inline

## ◆ offset()

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

CUTLASS_HOST_DEVICE LongIndex cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::offset | ( | TensorCoord const & | coord | ) | const
| inline

## ◆ operator+()

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

CUTLASS_HOST_DEVICE TensorRef cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::operator+ | ( | TensorCoord const & | b | ) | const
| inline

## ◆ operator+=()

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

CUTLASS_HOST_DEVICE TensorRef& cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::operator+= | ( | TensorCoord const & | b | ) |
| inline

## ◆ operator-()

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

CUTLASS_HOST_DEVICE TensorRef cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::operator- | ( | TensorCoord const & | b | ) | const
| inline

## ◆ operator-=()

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

CUTLASS_HOST_DEVICE TensorRef& cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::operator-= | ( | TensorCoord const & | b | ) |
| inline

## ◆ operator[]() [1/2]

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

CUTLASS_HOST_DEVICE Storage& cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::operator[] | ( | TensorCoord const & | coord | ) | const
| inline

## ◆ operator[]() [2/2]

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

CUTLASS_HOST_DEVICE Storage& cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::operator[] | ( | LongIndex | idx | ) | const
| inline

## ◆ reset() [1/2]

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

CUTLASS_HOST_DEVICE void cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::reset | ( | Storage * | ptr = `nullptr` | ) |
| inline

## ◆ reset() [2/2]

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

CUTLASS_HOST_DEVICE void cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::reset | ( | Storage * | ptr,
|  | StorageCoord const & | stride
| ) |  |
| inline

## ◆ stride() [1/2]

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

CUTLASS_HOST_DEVICE StorageCoord cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::stride | ( |  | ) | const
| inline

## ◆ stride() [2/2]

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

CUTLASS_HOST_DEVICE Index cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::stride | ( | int | dim | ) | const
| inline

## Member Data Documentation

## ◆ kRank

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

int const cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::kRank = Rank_
| static

## ◆ kStorageRank

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

int const cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::kStorageRank = StorageRank_
| static

## ◆ Rank

template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>

int const cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::Rank = kRank
| static

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

- tensor_ref.h