Cutlass CUDA Templates for Linear Algebra Subroutines and Solvers

Classes | Public Types | Public Member Functions | Static Public Member Functions | Static Public Attributes | List of all members
cutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ > Struct Template Reference

Iterator for accessing a stripmined tile in memory.

`#include <tile_iterator.h>`

Inheritance diagram for cutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >:

## Classes
struct | Params
| Parameters to the iterator. More...

## Public Types
typedef Traits_ | Traits
| concept TileTraits More...
typedef Scalar_ | Scalar
| Scalar element. More...
typedef FragmentElement_ | FragmentElement
| Fragment element. More...
typedef Index_ | Index
| Index type. More...
typedef long long | LongIndex
| Long index. More...
typedef Skew_ | Skew
| Skew quantity. More...
typedef Traits::Tile | Tile
| Tile shape. More...
typedef Traits::Delta | Delta
| Distance along each dimension. More...
typedef Traits::ImmediateOffsetStrides | ImmediateOffsetStrides
| The strides in each dimension between different loads/stores. More...
typedef Traits::Iterations | Iterations
| Iterations. More...
typedef Traits::ThreadOffset | ThreadOffset
| Thread offset. More...
typedef Vectorize< FragmentElement, kAccessSize >::Type | AccessType
| The elements loaded/store by one instruction. More...
typedef Fragment< Scalar, ShapeCount< Tile >::kCount, kFragmentSize > | Storage
| The storage. More...
typedef Fragment< FragmentElement, ShapeCount< Iterations >::kCount *kAccessSize > | Fragment
| The fragment. More...
typedef FragmentIterator< Fragment, Iterations, AccessType > | FragmentIterator
| The fragment iterator. More...
typedef FragmentConstIterator< Fragment, Iterations, AccessType > | FragmentConstIterator
| The fragment const iterator. More...
typedef FragmentIterator::FragmentShape | FragmentShape
| The shape of the fragment. More...
typedef PredicateVector< ShapeCount< Iterations >::kCount > | PredicateVector
| Default predicate mask type. More...

## Public Member Functions
CUTLASS_HOST_DEVICE bool | valid (int d, int h, int w, int c) const
| Is the iterator valid? More...

## Static Public Member Functions
template<typename PredicateIterator , typename PredicateFunctor >
static CUTLASS_HOST_DEVICE void | initialize_predicates (PredicateIterator predicate_it, PredicateFunctor const &predicate_func, Coord< 3 > const &offset)
| Initializes a predicate vector. More...

## Static Public Attributes
static IteratorAdvance::Kind const | kAdvance = Advance_
| Specifies dimension in which post-increment accesses advance. More...
static FragmentElementType::Kind const | kFragmentElementType = FragmentElementType_
| Specifies iterator storage fragment type (Scalar or WmmaMatrix) More...
static MemorySpace::Kind const | kMemorySpace = MemorySpace
| Source or destination memory space. More...
static int const | kAccessSize = Traits::kAccessSize
| The number of scalars accessed per load/store. More...
static int const | kFragmentSize
| The size of storage needed per fragment. More...

## Member Typedef Documentation

## ◆ AccessType

template<typename Traits_, typename Scalar_, IteratorAdvance::Kind Advance_ = IteratorAdvance::kH, MemorySpace::Kind MemorySpace = MemorySpace::kGeneric, typename Index_ = int, typename FragmentElement_ = Scalar_, FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar, typename Skew_ = Shape<0, 0, 0, 0>>

typedef Vectorize<FragmentElement, kAccessSize>::Type cutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >::AccessType

## ◆ Delta

template<typename Traits_, typename Scalar_, IteratorAdvance::Kind Advance_ = IteratorAdvance::kH, MemorySpace::Kind MemorySpace = MemorySpace::kGeneric, typename Index_ = int, typename FragmentElement_ = Scalar_, FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar, typename Skew_ = Shape<0, 0, 0, 0>>

typedef Traits::Delta cutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >::Delta

## ◆ Fragment

template<typename Traits_, typename Scalar_, IteratorAdvance::Kind Advance_ = IteratorAdvance::kH, MemorySpace::Kind MemorySpace = MemorySpace::kGeneric, typename Index_ = int, typename FragmentElement_ = Scalar_, FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar, typename Skew_ = Shape<0, 0, 0, 0>>

typedef Fragment<FragmentElement, ShapeCount<Iterations>::kCount * kAccessSize> cutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >::Fragment

## ◆ FragmentConstIterator

template<typename Traits_, typename Scalar_, IteratorAdvance::Kind Advance_ = IteratorAdvance::kH, MemorySpace::Kind MemorySpace = MemorySpace::kGeneric, typename Index_ = int, typename FragmentElement_ = Scalar_, FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar, typename Skew_ = Shape<0, 0, 0, 0>>

typedef FragmentConstIterator<Fragment, Iterations, AccessType> cutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >::FragmentConstIterator

## ◆ FragmentElement

template<typename Traits_, typename Scalar_, IteratorAdvance::Kind Advance_ = IteratorAdvance::kH, MemorySpace::Kind MemorySpace = MemorySpace::kGeneric, typename Index_ = int, typename FragmentElement_ = Scalar_, FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar, typename Skew_ = Shape<0, 0, 0, 0>>

typedef FragmentElement_ cutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >::FragmentElement

## ◆ FragmentIterator

template<typename Traits_, typename Scalar_, IteratorAdvance::Kind Advance_ = IteratorAdvance::kH, MemorySpace::Kind MemorySpace = MemorySpace::kGeneric, typename Index_ = int, typename FragmentElement_ = Scalar_, FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar, typename Skew_ = Shape<0, 0, 0, 0>>

typedef FragmentIterator<Fragment, Iterations, AccessType> cutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >::FragmentIterator

## ◆ FragmentShape

template<typename Traits_, typename Scalar_, IteratorAdvance::Kind Advance_ = IteratorAdvance::kH, MemorySpace::Kind MemorySpace = MemorySpace::kGeneric, typename Index_ = int, typename FragmentElement_ = Scalar_, FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar, typename Skew_ = Shape<0, 0, 0, 0>>

typedef FragmentIterator::FragmentShape cutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >::FragmentShape

## ◆ ImmediateOffsetStrides

template<typename Traits_, typename Scalar_, IteratorAdvance::Kind Advance_ = IteratorAdvance::kH, MemorySpace::Kind MemorySpace = MemorySpace::kGeneric, typename Index_ = int, typename FragmentElement_ = Scalar_, FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar, typename Skew_ = Shape<0, 0, 0, 0>>

typedef Traits::ImmediateOffsetStrides cutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >::ImmediateOffsetStrides

## ◆ Index

template<typename Traits_, typename Scalar_, IteratorAdvance::Kind Advance_ = IteratorAdvance::kH, MemorySpace::Kind MemorySpace = MemorySpace::kGeneric, typename Index_ = int, typename FragmentElement_ = Scalar_, FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar, typename Skew_ = Shape<0, 0, 0, 0>>

typedef Index_ cutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >::Index

## ◆ Iterations

template<typename Traits_, typename Scalar_, IteratorAdvance::Kind Advance_ = IteratorAdvance::kH, MemorySpace::Kind MemorySpace = MemorySpace::kGeneric, typename Index_ = int, typename FragmentElement_ = Scalar_, FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar, typename Skew_ = Shape<0, 0, 0, 0>>

typedef Traits::Iterations cutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >::Iterations

## ◆ LongIndex

template<typename Traits_, typename Scalar_, IteratorAdvance::Kind Advance_ = IteratorAdvance::kH, MemorySpace::Kind MemorySpace = MemorySpace::kGeneric, typename Index_ = int, typename FragmentElement_ = Scalar_, FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar, typename Skew_ = Shape<0, 0, 0, 0>>

typedef long long cutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >::LongIndex

## ◆ PredicateVector

template<typename Traits_, typename Scalar_, IteratorAdvance::Kind Advance_ = IteratorAdvance::kH, MemorySpace::Kind MemorySpace = MemorySpace::kGeneric, typename Index_ = int, typename FragmentElement_ = Scalar_, FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar, typename Skew_ = Shape<0, 0, 0, 0>>

typedef PredicateVector<ShapeCount<Iterations>::kCount> cutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >::PredicateVector

## ◆ Scalar

template<typename Traits_, typename Scalar_, IteratorAdvance::Kind Advance_ = IteratorAdvance::kH, MemorySpace::Kind MemorySpace = MemorySpace::kGeneric, typename Index_ = int, typename FragmentElement_ = Scalar_, FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar, typename Skew_ = Shape<0, 0, 0, 0>>

typedef Scalar_ cutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >::Scalar

## ◆ Skew

template<typename Traits_, typename Scalar_, IteratorAdvance::Kind Advance_ = IteratorAdvance::kH, MemorySpace::Kind MemorySpace = MemorySpace::kGeneric, typename Index_ = int, typename FragmentElement_ = Scalar_, FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar, typename Skew_ = Shape<0, 0, 0, 0>>

typedef Skew_ cutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >::Skew

## ◆ Storage

template<typename Traits_, typename Scalar_, IteratorAdvance::Kind Advance_ = IteratorAdvance::kH, MemorySpace::Kind MemorySpace = MemorySpace::kGeneric, typename Index_ = int, typename FragmentElement_ = Scalar_, FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar, typename Skew_ = Shape<0, 0, 0, 0>>

typedef Fragment<Scalar, ShapeCount<Tile>::kCount, kFragmentSize> cutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >::Storage

## ◆ ThreadOffset

template<typename Traits_, typename Scalar_, IteratorAdvance::Kind Advance_ = IteratorAdvance::kH, MemorySpace::Kind MemorySpace = MemorySpace::kGeneric, typename Index_ = int, typename FragmentElement_ = Scalar_, FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar, typename Skew_ = Shape<0, 0, 0, 0>>

typedef Traits::ThreadOffset cutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >::ThreadOffset

## ◆ Tile

template<typename Traits_, typename Scalar_, IteratorAdvance::Kind Advance_ = IteratorAdvance::kH, MemorySpace::Kind MemorySpace = MemorySpace::kGeneric, typename Index_ = int, typename FragmentElement_ = Scalar_, FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar, typename Skew_ = Shape<0, 0, 0, 0>>

typedef Traits::Tile cutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >::Tile

## ◆ Traits

template<typename Traits_, typename Scalar_, IteratorAdvance::Kind Advance_ = IteratorAdvance::kH, MemorySpace::Kind MemorySpace = MemorySpace::kGeneric, typename Index_ = int, typename FragmentElement_ = Scalar_, FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar, typename Skew_ = Shape<0, 0, 0, 0>>

typedef Traits_ cutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >::Traits

## Member Function Documentation

## ◆ initialize_predicates()

template<typename Traits_, typename Scalar_, IteratorAdvance::Kind Advance_ = IteratorAdvance::kH, MemorySpace::Kind MemorySpace = MemorySpace::kGeneric, typename Index_ = int, typename FragmentElement_ = Scalar_, FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar, typename Skew_ = Shape<0, 0, 0, 0>>
template<typename PredicateIterator , typename PredicateFunctor >

static CUTLASS_HOST_DEVICE void cutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >::initialize_predicates | ( | PredicateIterator | predicate_it,
|  | PredicateFunctor const & | predicate_func,
|  | Coord< 3 > const & | offset
| ) |  |
| inline static

## ◆ valid()

template<typename Traits_, typename Scalar_, IteratorAdvance::Kind Advance_ = IteratorAdvance::kH, MemorySpace::Kind MemorySpace = MemorySpace::kGeneric, typename Index_ = int, typename FragmentElement_ = Scalar_, FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar, typename Skew_ = Shape<0, 0, 0, 0>>

CUTLASS_HOST_DEVICE bool cutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >::valid | ( | int | d,
|  | int | h,
|  | int | w,
|  | int | c
| ) |  | const
| inline

## Member Data Documentation

## ◆ kAccessSize

template<typename Traits_, typename Scalar_, IteratorAdvance::Kind Advance_ = IteratorAdvance::kH, MemorySpace::Kind MemorySpace = MemorySpace::kGeneric, typename Index_ = int, typename FragmentElement_ = Scalar_, FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar, typename Skew_ = Shape<0, 0, 0, 0>>

int const cutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >::kAccessSize = Traits::kAccessSize
| static

## ◆ kAdvance

template<typename Traits_, typename Scalar_, IteratorAdvance::Kind Advance_ = IteratorAdvance::kH, MemorySpace::Kind MemorySpace = MemorySpace::kGeneric, typename Index_ = int, typename FragmentElement_ = Scalar_, FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar, typename Skew_ = Shape<0, 0, 0, 0>>

IteratorAdvance::Kind const cutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >::kAdvance = Advance_
| static

## ◆ kFragmentElementType

template<typename Traits_, typename Scalar_, IteratorAdvance::Kind Advance_ = IteratorAdvance::kH, MemorySpace::Kind MemorySpace = MemorySpace::kGeneric, typename Index_ = int, typename FragmentElement_ = Scalar_, FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar, typename Skew_ = Shape<0, 0, 0, 0>>

FragmentElementType::Kind const cutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >::kFragmentElementType = FragmentElementType_
| static

## ◆ kFragmentSize

template<typename Traits_, typename Scalar_, IteratorAdvance::Kind Advance_ = IteratorAdvance::kH, MemorySpace::Kind MemorySpace = MemorySpace::kGeneric, typename Index_ = int, typename FragmentElement_ = Scalar_, FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar, typename Skew_ = Shape<0, 0, 0, 0>>

int const cutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >::kFragmentSize
| static

Initial value:
=
(kFragmentElementType == FragmentElementType::kWmmaMatrix ? 16 : sizeof(AccessType))

## ◆ kMemorySpace

template<typename Traits_, typename Scalar_, IteratorAdvance::Kind Advance_ = IteratorAdvance::kH, MemorySpace::Kind MemorySpace = MemorySpace::kGeneric, typename Index_ = int, typename FragmentElement_ = Scalar_, FragmentElementType::Kind FragmentElementType_ = FragmentElementType::kScalar, typename Skew_ = Shape<0, 0, 0, 0>>

MemorySpace::Kind const cutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >::kMemorySpace = MemorySpace
| static

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

- tile_iterator.h