predicated_tile_access_iterator_2dthreadtile.h File Reference

Templates calculating the address and predicates to the load of tiles from pitch-linear rank=2 tensors. More...

#include "cutlass/array.h"
#include "cutlass/coord.h"
#include "cutlass/cutlass.h"
#include "cutlass/layout/matrix.h"
#include "cutlass/layout/pitch_linear.h"
#include "cutlass/matrix_shape.h"
#include "cutlass/predicate_vector.h"
#include "cutlass/tensor_ref.h"
#include "cutlass/tensor_view.h"

Include dependency graph for predicated_tile_access_iterator_2dthreadtile.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::PredicatedTileAccessIterator2dThreadTile< Shape, Element, Layout, AdvanceRank, ThreadMap, AccessType >
class	cutlass::transform::threadblock::PredicatedTileAccessIterator2dThreadTile< Shape_, Element_, layout::PitchLinear, AdvanceRank, ThreadMap_, AccessType_ >
class	cutlass::transform::threadblock::PredicatedTileAccessIterator2dThreadTile< Shape_, Element_, layout::PitchLinear, AdvanceRank, ThreadMap_, AccessType_ >::Params
	Parameters object is precomputed state and is host-constructible. More...
class	cutlass::transform::threadblock::PredicatedTileAccessIterator2dThreadTile< Shape_, Element_, layout::ColumnMajor, AdvanceRank, ThreadMap_, AccessType_ >
class	cutlass::transform::threadblock::PredicatedTileAccessIterator2dThreadTile< Shape_, Element_, layout::ColumnMajor, AdvanceRank, ThreadMap_, AccessType_ >::Params
	Parameters object is precomputed state and is host-constructible. More...
class	cutlass::transform::threadblock::PredicatedTileAccessIterator2dThreadTile< Shape_, Element_, layout::RowMajor, AdvanceRank, ThreadMap_, AccessType_ >
class	cutlass::transform::threadblock::PredicatedTileAccessIterator2dThreadTile< Shape_, Element_, layout::RowMajor, AdvanceRank, ThreadMap_, AccessType_ >::Params
	Parameters object is precomputed state and is host-constructible. More...

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.