CUTLASS
CUDA Templates for Linear Algebra Subroutines and Solvers
Public Types | Public Member Functions | Static Public Attributes | List of all members
cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment > Class Template Reference

#include <regular_tile_iterator_tensor_op.h>

Public Types

using Shape = Shape_
 
using Element = Element_
 
using Layout = layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element))>
 
using Index = typename Layout::Index
 
using LongIndex = typename Layout::LongIndex
 
using TensorRef = TensorRef< Element, Layout >
 
using TensorCoord = typename Layout::TensorCoord
 
using ThreadMap = ThreadMap_
 
using UnderlyingIterator = RegularTileIterator< layout::PitchLinearShape< Shape::kColumn, Shape::kRow >, Element, layout::TensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element))>,(kAdvanceRank==0?1:0), ThreadMap_ >
 Underlying iterator type. More...
 
using Fragment = Array< Element, UnderlyingIterator::Fragment::kElements >
 Fragment object to be loaded or stored. More...
 

Public Member Functions

CUTLASS_HOST_DEVICE RegularTileIterator (TensorRef ref, int thread_id)
 Construct a TileIterator with zero threadblock offset. More...
 
CUTLASS_HOST_DEVICE void add_pointer_offset (LongIndex pointer_offset)
 Adds a pointer offset in units of Element. More...
 
CUTLASS_DEVICE void add_tile_offset (TensorCoord const &coord)
 Adds a tile offset. More...
 
CUTLASS_HOST_DEVICE RegularTileIteratoroperator++ ()
 Advances to the next tile in memory. More...
 
CUTLASS_HOST_DEVICE RegularTileIterator operator++ (int)
 Advances to the next tile in memory. More...
 
CUTLASS_DEVICE void load_with_pointer_offset (Fragment &frag, Index pointer_offset)
 Loads a fragment from memory. More...
 
CUTLASS_DEVICE void load (Fragment &frag)
 Loads a fragment from memory. More...
 
CUTLASS_DEVICE void store_with_pointer_offset (Fragment const &frag, Index pointer_offset)
 Store a fragment to memory. More...
 
CUTLASS_DEVICE void store (Fragment const &frag)
 Store a fragment to memory. More...
 

Static Public Attributes

static int const kAdvanceRank = AdvanceRank
 
static int const kAlignment = Alignment
 

Detailed Description

template<typename Shape_, typename Element_, int AdvanceRank, typename ThreadMap_, int Alignment>
class cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >

Tile Iterator specialized for row-major congruous TensorOp formats.

Satisfies: ForwardTileIteratorConcept | ReadableContiguousTileIteratorConcept | WriteableContiguousTileIteratorConcept

Member Typedef Documentation

template<typename Shape_ , typename Element_ , int AdvanceRank, typename ThreadMap_ , int Alignment>
using cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::Element = Element_
template<typename Shape_ , typename Element_ , int AdvanceRank, typename ThreadMap_ , int Alignment>
using cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::Fragment = Array<Element, UnderlyingIterator::Fragment::kElements>
template<typename Shape_ , typename Element_ , int AdvanceRank, typename ThreadMap_ , int Alignment>
using cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::Index = typename Layout::Index
template<typename Shape_ , typename Element_ , int AdvanceRank, typename ThreadMap_ , int Alignment>
using cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::Layout = layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits<Element_>::value, int(128 / sizeof(Element))>
template<typename Shape_ , typename Element_ , int AdvanceRank, typename ThreadMap_ , int Alignment>
using cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::LongIndex = typename Layout::LongIndex
template<typename Shape_ , typename Element_ , int AdvanceRank, typename ThreadMap_ , int Alignment>
using cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::Shape = Shape_
template<typename Shape_ , typename Element_ , int AdvanceRank, typename ThreadMap_ , int Alignment>
using cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::TensorCoord = typename Layout::TensorCoord
template<typename Shape_ , typename Element_ , int AdvanceRank, typename ThreadMap_ , int Alignment>
using cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::TensorRef = TensorRef<Element, Layout>
template<typename Shape_ , typename Element_ , int AdvanceRank, typename ThreadMap_ , int Alignment>
using cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::ThreadMap = ThreadMap_
template<typename Shape_ , typename Element_ , int AdvanceRank, typename ThreadMap_ , int Alignment>
using cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::UnderlyingIterator = RegularTileIterator< layout::PitchLinearShape<Shape::kColumn, Shape::kRow>, Element, layout::TensorOpMultiplicandCongruous<sizeof_bits<Element_>::value, int(128 / sizeof(Element))>, (kAdvanceRank == 0 ? 1 : 0), ThreadMap_>

Constructor & Destructor Documentation

template<typename Shape_ , typename Element_ , int AdvanceRank, typename ThreadMap_ , int Alignment>
CUTLASS_HOST_DEVICE cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::RegularTileIterator ( TensorRef  ref,
int  thread_id 
)
inline
Parameters
refPointer to start of tensor
thread_idID of each participating thread

Member Function Documentation

template<typename Shape_ , typename Element_ , int AdvanceRank, typename ThreadMap_ , int Alignment>
CUTLASS_HOST_DEVICE void cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::add_pointer_offset ( LongIndex  pointer_offset)
inline
template<typename Shape_ , typename Element_ , int AdvanceRank, typename ThreadMap_ , int Alignment>
CUTLASS_DEVICE void cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::add_tile_offset ( TensorCoord const &  coord)
inline
template<typename Shape_ , typename Element_ , int AdvanceRank, typename ThreadMap_ , int Alignment>
CUTLASS_DEVICE void cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::load ( Fragment frag)
inline
template<typename Shape_ , typename Element_ , int AdvanceRank, typename ThreadMap_ , int Alignment>
CUTLASS_DEVICE void cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::load_with_pointer_offset ( Fragment frag,
Index  pointer_offset 
)
inline
template<typename Shape_ , typename Element_ , int AdvanceRank, typename ThreadMap_ , int Alignment>
CUTLASS_HOST_DEVICE RegularTileIterator& cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::operator++ ( )
inline
template<typename Shape_ , typename Element_ , int AdvanceRank, typename ThreadMap_ , int Alignment>
CUTLASS_HOST_DEVICE RegularTileIterator cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::operator++ ( int  )
inline
template<typename Shape_ , typename Element_ , int AdvanceRank, typename ThreadMap_ , int Alignment>
CUTLASS_DEVICE void cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::store ( Fragment const &  frag)
inline
template<typename Shape_ , typename Element_ , int AdvanceRank, typename ThreadMap_ , int Alignment>
CUTLASS_DEVICE void cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::store_with_pointer_offset ( Fragment const &  frag,
Index  pointer_offset 
)
inline

Member Data Documentation

template<typename Shape_ , typename Element_ , int AdvanceRank, typename ThreadMap_ , int Alignment>
int const cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::kAdvanceRank = AdvanceRank
static
template<typename Shape_ , typename Element_ , int AdvanceRank, typename ThreadMap_ , int Alignment>
int const cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value, int(128/sizeof(Element_))>, AdvanceRank, ThreadMap_, Alignment >::kAlignment = Alignment
static

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