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

#include <igemm_global_tile.h>

Inheritance diagram for cutlass::gemm::IgemmGlobalIteratorAb< TileTraits_, Index_ >:
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 > >

Public Types

typedef GemmGlobalIteratorAb< TileTraits_, Index_ > Base
 The base class. More...
 
typedef TileTraits_::ThreadOffset ThreadOffset
 The functor to compute the thread offset. More...
 
- Public Types inherited from cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >
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_DEVICE IgemmGlobalIteratorAb (typename Base::Params const &_params, const Coord< 3 > &threadblock_offset, ThreadOffset thread_offset_func=ThreadOffset())
 Constructor. More...
 
CUTLASS_DEVICE void initialize_predicates (const Coord< 3 > &bounds, const Coord< 3 > &threadblock_offset)
 
CUTLASS_DEVICE void load_element (typename Base::AccessType &value, int d, int h, int w, int c) const
 
- Public Member Functions inherited from cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >
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

uint32_t mask_
 The mask to clean up the values. More...
 
- Public Attributes inherited from cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >
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...
 

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...
 
- Static Public Attributes inherited from cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >
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...
 

Member Typedef Documentation

◆ Base

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

◆ ThreadOffset

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

Constructor & Destructor Documentation

◆ IgemmGlobalIteratorAb()

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

Member Function Documentation

◆ initialize_predicates()

template<typename TileTraits_ , typename Index_ = int>
CUTLASS_DEVICE void cutlass::gemm::IgemmGlobalIteratorAb< TileTraits_, Index_ >::initialize_predicates ( const Coord< 3 > &  bounds,
const Coord< 3 > &  threadblock_offset 
)
inline

◆ load_element()

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

Member Data Documentation

◆ mask_

template<typename TileTraits_ , typename Index_ = int>
uint32_t cutlass::gemm::IgemmGlobalIteratorAb< TileTraits_, Index_ >::mask_

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