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.