regular_tile_iterator_tensor_op_sm70.h File Reference

Templates implementing loading of tiles from pitch-linear rank=2 tensors. More...

`#include "cutlass/cutlass.h"`
`#include "cutlass/array.h"`
`#include "cutlass/matrix_coord.h"`
`#include "cutlass/tensor_ref.h"`
`#include "cutlass/layout/pitch_linear.h"`
`#include "cutlass/layout/tensor_op_multiplicand_sm70.h"`
`#include "cutlass/transform/threadblock/regular_tile_iterator.h"`
Include dependency graph for regular_tile_iterator_tensor_op_sm70.h:
This graph shows which files directly or indirectly include this file:

Go to the source code of this file.

## Classes

class | cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, AdvanceRank, ThreadMap_, Alignment >
struct | cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, AdvanceRank, ThreadMap_, Alignment >::Detail
| Internal details made public to facilitate introspection. More...
class | cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::ColumnMajorVoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, AdvanceRank, ThreadMap_, Alignment >
class | cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::RowMajorVoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, AdvanceRank, ThreadMap_, Alignment >
class | cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::VoltaTensorOpMultiplicandBCongruous< sizeof_bits< Element_ >::value >, AdvanceRank, ThreadMap_, Alignment >
struct | cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::VoltaTensorOpMultiplicandBCongruous< sizeof_bits< Element_ >::value >, AdvanceRank, ThreadMap_, Alignment >::Detail
| Internal details made public to facilitate introspection. More...
class | cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::ColumnMajorVoltaTensorOpMultiplicandBCongruous< sizeof_bits< Element_ >::value >, AdvanceRank, ThreadMap_, Alignment >
class | cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::RowMajorVoltaTensorOpMultiplicandBCongruous< sizeof_bits< Element_ >::value >, AdvanceRank, ThreadMap_, Alignment >
class | cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::VoltaTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Shape_::kContiguous >, AdvanceRank, ThreadMap_, Alignment >
struct | cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::VoltaTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Shape_::kContiguous >, AdvanceRank, ThreadMap_, Alignment >::Detail
| Internal details made public to facilitate introspection. More...
class | cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::ColumnMajorVoltaTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Shape_::kRow >, AdvanceRank, ThreadMap_, Alignment >
class | cutlass::transform::threadblock::RegularTileIterator< Shape_, Element_, layout::RowMajorVoltaTensorOpMultiplicandCrosswise< sizeof_bits< Element_ >::value, Shape_::kColumn >, AdvanceRank, ThreadMap_, Alignment >

## Namespaces

| cutlass
| cutlass::transform
| cutlass::transform::threadblock

## Detailed Description

This iterator uses masks to guard out-of-bounds accesses and visits the last "residue" tile first, with the objective of minimizing predicate mask updates during steady-state operation.

A precomputed "Params" object minimizes the amount of state that must be stored in registers, and integer addition is used to advance the pointer through memory.