Cutlass
CUDA Templates for Linear Algebra Subroutines and Solvers
Classes | Public Types | Public Member Functions | Public Attributes | Static Public Attributes | List of all members
cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ > Struct Template Reference

#include <gemm_global_tile.h>

Inheritance diagram for cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >:
cutlass::TileLoadIterator< TileTraits_, TileTraits_::Scalar, TileTraits_::MultiplicandTraits::kKstrided ? IteratorAdvance::kH :IteratorAdvance::kW, MemorySpace::kGlobal, Index_ > cutlass::TileIteratorBase< TileTraits_, TileTraits_::Scalar, Advance_, MemorySpace, Index_, TileTraits_::Scalar, FragmentElementType::kScalar, Shape< 0, 0, 0, 0 > > cutlass::gemm::IgemmGlobalIteratorAb< TileTraits_, Index_ >

Classes

struct  Params
 

Public Types

typedef GemmGlobalIteratorAb< TileTraits_, Index_ > This_
 This class. More...
 
typedef TileLoadIterator< TileTraits_, typename TileTraits_::Scalar, TileTraits_::MultiplicandTraits::kKstrided ? IteratorAdvance::kH :IteratorAdvance::kW, MemorySpace::kGlobal, Index_ > Base
 The base class. More...
 
typedef TileTraits_::Tile Tile
 The tile. More...
 
typedef Base::Fragment Fragment
 Fragment type loaded by the iterator. More...
 
typedef TileTraits_::Scalar Scalar
 The scalar. More...
 
typedef TileTraits_::Threads Threads
 The threads. More...
 
typedef Index_ Index
 The index. More...
 
typedef long long LongIndex
 Long index. More...
 
typedef TileTraits_::ThreadOffset ThreadOffset
 The thread offset. More...
 
typedef cutlass::PredicateVector< ShapeCount< typename Base::Iterations >::kCount > PredicateVector
 
typedef Base::Params BaseParams
 Iterator parameters type. More...
 
- Public Types inherited from cutlass::TileLoadIterator< TileTraits_, TileTraits_::Scalar, TileTraits_::MultiplicandTraits::kKstrided ? IteratorAdvance::kH :IteratorAdvance::kW, MemorySpace::kGlobal, Index_ >
enum  
 Do we require a fence? More...
 
typedef TileIteratorBase< TileTraits_, TileTraits_::Scalar, Advance_, MemorySpace, Index_, TileTraits_::Scalar, FragmentElementType::kScalar, Shape< 0, 0, 0, 0 > > Base
 Base class. More...
 
typedef Base::Traits Traits
 concept TileTraits More...
 
typedef Base::Scalar Scalar
 Scalar element. More...
 
typedef TileTraits_::Scalar FragmentElement
 Fragment element. More...
 
typedef Base::Index Index
 Index type. More...
 
typedef Base::LongIndex LongIndex
 Index type. More...
 
typedef Base::Skew Skew
 Skew quantity. More...
 
typedef Base::Tile Tile
 Tile shape. More...
 
typedef Base::Delta Delta
 Delta. More...
 
typedef Base::Iterations Iterations
 Iterations. More...
 
typedef Base::ThreadOffset ThreadOffset
 ThreadOffset functor. More...
 
typedef Base::FragmentShape FragmentShape
 Fragment type. More...
 
typedef Base::AccessType AccessType
 Memory access type. More...
 
typedef Base::Fragment Fragment
 Fragment definition. More...
 
typedef Base::FragmentIterator FragmentIterator
 Fragment iterator definition. More...
 
typedef Base::FragmentConstIterator FragmentConstIterator
 Fragment const iterator definition. More...
 
typedef Base::PredicateVector PredicateVector
 Default predicate mask type. More...
 
typedef Base::Storage SharedStorage
 Storage object that may be loaded from. More...
 
typedef Base::Params BaseParams
 IteratorBase parameters. More...
 
typedef Scalar const * Pointer
 The pointer type. More...
 
typedef TensorRef< Scalar const, 4 > TensorRef
 Tensor reference for the load iterator. More...
 
- Public Types inherited from cutlass::TileIteratorBase< TileTraits_, TileTraits_::Scalar, Advance_, MemorySpace, Index_, TileTraits_::Scalar, FragmentElementType::kScalar, Shape< 0, 0, 0, 0 > >
typedef TileTraits_ Traits
 concept TileTraits More...
 
typedef TileTraits_::Scalar Scalar
 Scalar element. More...
 
typedef TileTraits_::Scalar FragmentElement
 Fragment element. More...
 
typedef Index_ Index
 Index type. More...
 
typedef long long LongIndex
 Long index. More...
 
typedef Shape< 0, 0, 0, 0 > 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 void initialize_predicates (const Coord< 3 > &bounds, const Coord< 3 > &block_offset)
 
CUTLASS_HOST_DEVICE GemmGlobalIteratorAb (Params const &_params, const Coord< 3 > &threadblock_offset, ThreadOffset thread_offset_func=ThreadOffset())
 Ctor. More...
 
CUTLASS_HOST_DEVICE void inc_w ()
 Increment the pointer in the W dimension. More...
 
CUTLASS_HOST_DEVICE void inc_h ()
 Increment the pointer in the H dimension. More...
 
CUTLASS_HOST_DEVICE void inc_d ()
 Increment the pointer in the D dimension. More...
 
CUTLASS_HOST_DEVICE void inc_advance ()
 Increment the pointer to move to the next iteration. More...
 
CUTLASS_HOST_DEVICE void load_element (typename Base::AccessType &value, int d, int h, int w, int c) const
 Loads a single fragment element from memory. More...
 
CUTLASS_HOST_DEVICE void residue (Index k)
 That's the residue! Update the predicates. More...
 
CUTLASS_HOST_DEVICE bool valid (int d, int h, int w, int c) const
 Is the valid? More...
 
CUTLASS_HOST_DEVICE GemmGlobalIteratorAboperator+= (Coord< 3 > const &offset)
 Adds a vector offset to the iterator. More...
 
CUTLASS_HOST_DEVICE void add_pointer_offset (Index offset)
 
CUTLASS_HOST_DEVICE Index stride_advance (void)
 
template<typename Fragment >
CUTLASS_HOST_DEVICE void load_post_increment (Fragment &fragment)
 
- Public Member Functions inherited from cutlass::TileLoadIterator< TileTraits_, TileTraits_::Scalar, TileTraits_::MultiplicandTraits::kKstrided ? IteratorAdvance::kH :IteratorAdvance::kW, MemorySpace::kGlobal, Index_ >
CUTLASS_HOST_DEVICE void initialize_predicates (PredicateIterator predicate_it, Coord< 3 > const &bounds, Coord< 3 > const &block_offset=make_Coord(0, 0, 0))
 Initializes a predicate vector using a RegularTilePredicateFunctor. More...
 
CUTLASS_HOST_DEVICE void initialize_predicates (PredicateIterator predicate_it, PredicateFunctor const &functor, Coord< 3 > const &block_offset)
 Initializes a predicate vector using an arbitrary predicate functor. More...
 
CUTLASS_HOST_DEVICE TileLoadIterator ()
 Default constructor. More...
 
CUTLASS_HOST_DEVICE TileLoadIterator (Params const &_params, Coord< 3 > const &block_offset=make_Coord(0, 0, 0), ThreadOffset thread_offset_func=ThreadOffset())
 Constructs a tile load iterator. More...
 
CUTLASS_HOST_DEVICE TileLoadIterator (Params const &, Scalar const *ptr, Coord< 3 > const &block_offset=make_Coord(0, 0, 0), ThreadOffset thread_offset_func=ThreadOffset())
 Constructs a tile load iterator. More...
 
CUTLASS_HOST_DEVICE void inc_d ()
 Increment in the D dimension. More...
 
CUTLASS_HOST_DEVICE void inc_h ()
 Increment in the H dimension. More...
 
CUTLASS_HOST_DEVICE void inc_w ()
 Increment in the W dimension. More...
 
CUTLASS_HOST_DEVICE void inc_advance ()
 Increment in the next dimension. More...
 
CUTLASS_HOST_DEVICE void load_element (AccessType &value, int d, int h, int w, int c) const
 Loads a single fragment element from memory. More...
 
CUTLASS_HOST_DEVICE void inc_stage ()
 Increment the stage. More...
 
CUTLASS_HOST_DEVICE TileLoadIteratoroperator+= (Coord< 3 > const &offset)
 Adds a vector offset to the iterator. More...
 
CUTLASS_HOST_DEVICE void add_pointer_offset (LongIndex offset)
 Adds a raw offset to the pointer. More...
 
CUTLASS_HOST_DEVICE Index stride_advance (void)
 
CUTLASS_HOST_DEVICE void load_post_increment (Fragment &fragment, PredicateIterator pred_it)
 Loads a fragment and advances the iterator to the next tile. More...
 
CUTLASS_HOST_DEVICE void load_post_increment (Fragment &fragment)
 Loads a fragment and advances the iterator to the next tile. More...
 
CUTLASS_HOST_DEVICE void load (Fragment &fragment, PredicateIterator pred_it) const
 Loads a fragment without advancing the iterator.. More...
 
CUTLASS_HOST_DEVICE void load (Fragment &fragment) const
 Loads a fragment without advancing the iterator.. More...
 
CUTLASS_HOST_DEVICE void load (Fragment &fragment, int d)
 Loads a fragment without advancing the iterator.. More...
 
- Public Member Functions inherited from cutlass::TileIteratorBase< TileTraits_, TileTraits_::Scalar, Advance_, MemorySpace, Index_, TileTraits_::Scalar, FragmentElementType::kScalar, Shape< 0, 0, 0, 0 > >
CUTLASS_HOST_DEVICE bool valid (int d, int h, int w, int c) const
 Is the iterator valid? More...
 

Public Attributes

Coord< 4 > thread_offset
 Offset of an individual lane from the start of the tile. More...
 
Params params
 The parameters. More...
 
PredicateVector predicates
 The predicates. More...
 
- Public Attributes inherited from cutlass::TileLoadIterator< TileTraits_, TileTraits_::Scalar, TileTraits_::MultiplicandTraits::kKstrided ? IteratorAdvance::kH :IteratorAdvance::kW, MemorySpace::kGlobal, Index_ >
Params params
 Parameters structure. More...
 
Coord< 4 > thread_offset
 Offset of an individual lane from the start of the tile. More...
 
int stage
 Stage argument enables wrapping after some number of tiles have been loaded. More...
 

Static Public Attributes

static MatrixLayout::Kind const kLayout = TileTraits_::kLayout
 The layout. More...
 
static IteratorAdvance::Kind const kAdvance = Base::kAdvance
 Specifies in which dimension post-increment accesses advance. More...
 
- Static Public Attributes inherited from cutlass::TileLoadIterator< TileTraits_, TileTraits_::Scalar, TileTraits_::MultiplicandTraits::kKstrided ? IteratorAdvance::kH :IteratorAdvance::kW, MemorySpace::kGlobal, Index_ >
static IteratorAdvance::Kind const kAdvance
 Specifies in which dimension post-increment accesses advance. More...
 
static FragmentElementType::Kind const kFragmentElementType
 Specifies type of iterator fragment storage (Salar or WmmaMatrix) More...
 
static MemorySpace::Kind const kMemorySpace
 Source or destination memory space. More...
 
static int const kAccessSize
 The number of scalars accessed per load/store. More...
 
- Static Public Attributes inherited from cutlass::TileIteratorBase< TileTraits_, TileTraits_::Scalar, Advance_, MemorySpace, Index_, TileTraits_::Scalar, FragmentElementType::kScalar, Shape< 0, 0, 0, 0 > >
static IteratorAdvance::Kind const kAdvance
 Specifies dimension in which post-increment accesses advance. More...
 
static FragmentElementType::Kind const kFragmentElementType
 Specifies iterator storage fragment type (Scalar or WmmaMatrix) More...
 
static MemorySpace::Kind const kMemorySpace
 Source or destination memory space. More...
 
static int const kAccessSize
 The number of scalars accessed per load/store. More...
 
static int const kFragmentSize
 The size of storage needed per fragment. More...
 

Additional Inherited Members

- Static Public Member Functions inherited from cutlass::TileIteratorBase< TileTraits_, TileTraits_::Scalar, Advance_, MemorySpace, Index_, TileTraits_::Scalar, FragmentElementType::kScalar, Shape< 0, 0, 0, 0 > >
static CUTLASS_HOST_DEVICE void initialize_predicates (PredicateIterator predicate_it, PredicateFunctor const &predicate_func, Coord< 3 > const &offset)
 Initializes a predicate vector. More...
 

Member Typedef Documentation

◆ Base

template<typename TileTraits_ , typename Index_ = int>
typedef TileLoadIterator<TileTraits_, typename TileTraits_::Scalar, TileTraits_::MultiplicandTraits::kKstrided ? IteratorAdvance::kH : IteratorAdvance::kW, MemorySpace::kGlobal, Index_> cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::Base

◆ BaseParams

template<typename TileTraits_ , typename Index_ = int>
typedef Base::Params cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::BaseParams

◆ Fragment

template<typename TileTraits_ , typename Index_ = int>
typedef Base::Fragment cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::Fragment

◆ Index

template<typename TileTraits_ , typename Index_ = int>
typedef Index_ cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::Index

◆ LongIndex

template<typename TileTraits_ , typename Index_ = int>
typedef long long cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::LongIndex

◆ PredicateVector

template<typename TileTraits_ , typename Index_ = int>
typedef cutlass::PredicateVector<ShapeCount<typename Base::Iterations>::kCount> cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::PredicateVector

◆ Scalar

template<typename TileTraits_ , typename Index_ = int>
typedef TileTraits_::Scalar cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::Scalar

◆ This_

template<typename TileTraits_ , typename Index_ = int>
typedef GemmGlobalIteratorAb<TileTraits_, Index_> cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::This_

◆ ThreadOffset

template<typename TileTraits_ , typename Index_ = int>
typedef TileTraits_::ThreadOffset cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::ThreadOffset

◆ Threads

template<typename TileTraits_ , typename Index_ = int>
typedef TileTraits_::Threads cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::Threads

◆ Tile

template<typename TileTraits_ , typename Index_ = int>
typedef TileTraits_::Tile cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::Tile

Constructor & Destructor Documentation

◆ GemmGlobalIteratorAb()

template<typename TileTraits_ , typename Index_ = int>
CUTLASS_HOST_DEVICE cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::GemmGlobalIteratorAb ( Params const &  _params,
const Coord< 3 > &  threadblock_offset,
ThreadOffset  thread_offset_func = ThreadOffset() 
)
inline

Member Function Documentation

◆ add_pointer_offset()

template<typename TileTraits_ , typename Index_ = int>
CUTLASS_HOST_DEVICE void cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::add_pointer_offset ( Index  offset)
inline

◆ inc_advance()

template<typename TileTraits_ , typename Index_ = int>
CUTLASS_HOST_DEVICE void cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::inc_advance ( )
inline

◆ inc_d()

template<typename TileTraits_ , typename Index_ = int>
CUTLASS_HOST_DEVICE void cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::inc_d ( )
inline

◆ inc_h()

template<typename TileTraits_ , typename Index_ = int>
CUTLASS_HOST_DEVICE void cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::inc_h ( )
inline

◆ inc_w()

template<typename TileTraits_ , typename Index_ = int>
CUTLASS_HOST_DEVICE void cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::inc_w ( )
inline

◆ initialize_predicates()

template<typename TileTraits_ , typename Index_ = int>
CUTLASS_HOST_DEVICE void cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::initialize_predicates ( const Coord< 3 > &  bounds,
const Coord< 3 > &  block_offset 
)
inline

◆ load_element()

template<typename TileTraits_ , typename Index_ = int>
CUTLASS_HOST_DEVICE void cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::load_element ( typename Base::AccessType value,
int  d,
int  h,
int  w,
int  c 
) const
inline

◆ load_post_increment()

template<typename TileTraits_ , typename Index_ = int>
template<typename Fragment >
CUTLASS_HOST_DEVICE void cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::load_post_increment ( Fragment fragment)
inline

◆ operator+=()

template<typename TileTraits_ , typename Index_ = int>
CUTLASS_HOST_DEVICE GemmGlobalIteratorAb& cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::operator+= ( Coord< 3 > const &  offset)
inline

◆ residue()

template<typename TileTraits_ , typename Index_ = int>
CUTLASS_HOST_DEVICE void cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::residue ( Index  k)
inline

◆ stride_advance()

template<typename TileTraits_ , typename Index_ = int>
CUTLASS_HOST_DEVICE Index cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::stride_advance ( void  )
inline

◆ valid()

template<typename TileTraits_ , typename Index_ = int>
CUTLASS_HOST_DEVICE bool cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::valid ( int  d,
int  h,
int  w,
int  c 
) const
inline

Member Data Documentation

◆ kAdvance

template<typename TileTraits_ , typename Index_ = int>
IteratorAdvance::Kind const cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::kAdvance = Base::kAdvance
static

◆ kLayout

template<typename TileTraits_ , typename Index_ = int>
MatrixLayout::Kind const cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::kLayout = TileTraits_::kLayout
static

◆ params

template<typename TileTraits_ , typename Index_ = int>
Params cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::params

◆ predicates

template<typename TileTraits_ , typename Index_ = int>
PredicateVector cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::predicates

◆ thread_offset

template<typename TileTraits_ , typename Index_ = int>
Coord<4> cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >::thread_offset

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