| CUTLASS CUDA Templates for Linear Algebra Subroutines and Solvers

- Main Page
- Modules
- Namespaces
- Classes
- Files

- Class List
- Class Index
- Class Hierarchy
- Class Members

Classes | Public Types | Public Member Functions | Static Public Attributes | List of all members
cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::VoltaTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Shape_::kContiguous >, AdvanceRank, ThreadMap_, Alignment > Class Template Reference

`#include <regular_tile_iterator_tensor_op_sm70.h>`

## Classes
struct | Detail
| Internal details made public to facilitate introspection. More...

## Public Types
using | Shape = Shape_
using | Element = Element_
using | Layout = layout::VoltaTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Shape::kContiguous >
using | Index = typename Layout::Index
using | LongIndex = typename Layout::LongIndex
using | TensorRef = TensorRef< Element, Layout >
using | TensorCoord = typename Layout::TensorCoord
using | ThreadMap = ThreadMap_
using | Fragment = Array< Element, ThreadMap::Iterations::kCount *ThreadMap::kElementsPerAccess >
| 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_HOST_DEVICE RegularTileIterator & | operator++ ()
| 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 | add_tile_offset (TensorCoord const &coord)
| Adds a tile offset. 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

## Detailed Description

### template<typename Shape_, typename Element_, int AdvanceRank, typename ThreadMap_, int Alignment>
class cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::VoltaTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Shape_::kContiguous >, AdvanceRank, ThreadMap_, Alignment >

Tile iterator specialized for crosswise arrangements for TensorOps.

Volta TN SMEM layout is a little diffrent: Crosseised elements will be stored in a line, while contiguous elements sre stored in line-by-line. Padding is used to reduce SMEM bank conflicts.

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::VoltaTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Shape_::kContiguous >, AdvanceRank, ThreadMap_, Alignment >::Element = Element_

template<typename Shape_ , typename Element_ , int AdvanceRank, typename ThreadMap_ , int Alignment>

using cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::VoltaTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Shape_::kContiguous >, AdvanceRank, ThreadMap_, Alignment >::Fragment = Array<Element, ThreadMap::Iterations::kCount * ThreadMap::kElementsPerAccess>

template<typename Shape_ , typename Element_ , int AdvanceRank, typename ThreadMap_ , int Alignment>

using cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::VoltaTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Shape_::kContiguous >, 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::VoltaTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Shape_::kContiguous >, AdvanceRank, ThreadMap_, Alignment >::Layout = layout::VoltaTensorOpMultiplicandCrosswise<sizeof_bits<Element_>::value, Shape::kContiguous>

template<typename Shape_ , typename Element_ , int AdvanceRank, typename ThreadMap_ , int Alignment>

using cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::VoltaTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Shape_::kContiguous >, 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::VoltaTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Shape_::kContiguous >, AdvanceRank, ThreadMap_, Alignment >::Shape = Shape_

template<typename Shape_ , typename Element_ , int AdvanceRank, typename ThreadMap_ , int Alignment>

using cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::VoltaTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Shape_::kContiguous >, 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::VoltaTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Shape_::kContiguous >, 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::VoltaTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Shape_::kContiguous >, AdvanceRank, ThreadMap_, Alignment >::ThreadMap = 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::VoltaTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Shape_::kContiguous >, AdvanceRank, ThreadMap_, Alignment >::RegularTileIterator | ( | TensorRef | ref,
|  | int | thread_id
| ) |  |
| inline

Parameters

  ref | Pointer to start of tensor
  thread_id | ID 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::VoltaTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Shape_::kContiguous >, 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::VoltaTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Shape_::kContiguous >, 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::VoltaTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Shape_::kContiguous >, 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::VoltaTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Shape_::kContiguous >, 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::VoltaTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Shape_::kContiguous >, 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::VoltaTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Shape_::kContiguous >, 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::VoltaTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Shape_::kContiguous >, 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::VoltaTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Shape_::kContiguous >, 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::VoltaTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Shape_::kContiguous >, AdvanceRank, ThreadMap_, Alignment >::kAdvanceRank = AdvanceRank
| static

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

- regular_tile_iterator_tensor_op_sm70.h