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_ >:
cutlass::TileLoadIterator< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ > cutlass::TileStoreIterator< 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, kFragmentSizeStorage
 The storage. More...
 
typedef Fragment< FragmentElement, ShapeCount< Iterations >::kCount *kAccessSizeFragment
 The fragment. More...
 
typedef FragmentIterator< Fragment, Iterations, AccessTypeFragmentIterator
 The fragment iterator. More...
 
typedef FragmentConstIterator< Fragment, Iterations, AccessTypeFragmentConstIterator
 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 
)
inlinestatic

◆ 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

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