cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize > Class Template Reference

`#include <mma_simt_tile_iterator.h>`

## Public Types

using | Shape = Shape_
| Shape of tile to load (concept: MatrixShape) More...
using | Element = Element_
| Element type. More...
using | Layout = layout::ColumnMajorInterleaved< 4 >
| Layout of policy. More...
using | Policy = Policy_
| Decomposition of elements among threads. More...
using | TensorRef = TensorRef< Element, Layout >
| TensorRef type for loading element from a tensor. More...
using | Index = typename TensorRef::Index
| Index type. More...
using | LongIndex = typename TensorRef::LongIndex
| Long Index type. More...
using | TensorCoord = typename TensorRef::TensorCoord
| Coordinate for an element in the tensor. More...
using | ThreadShape = MatrixShape< Shape::kRow/Policy::WarpShape::kRow, Shape::kColumn >
| Thread-level shape of a fragment. More...
using | Iterations = MatrixShape< ThreadShape::kRow/Policy::LaneMmaShape::kM, ThreadShape::kColumn/Policy::LaneMmaShape::kK >
| Number of individual loads. More...
using | Fragment = Array< Element, ThreadShape::kCount >
| Fragment object holding a thread's part of a tile. More...

## Public Member Functions

CUTLASS_HOST_DEVICE | MmaSimtTileIterator ()
CUTLASS_HOST_DEVICE | MmaSimtTileIterator (TensorRef ref, int lane_id)
| Constructor from TensorRef. More...
CUTLASS_HOST_DEVICE MmaSimtTileIterator & | add_pointer_offset (LongIndex offset)
| Adds a pointer offset to internal pointer(s) to advance through memory. More...
CUTLASS_HOST_DEVICE MmaSimtTileIterator & | add_tile_offset (TensorCoord const &coord)
| Advances an iterator along logical dimensions of matrix in units of whole tiles. More...
CUTLASS_HOST_DEVICE MmaSimtTileIterator & | operator++ ()
| Advances the iterator along the advance dimension. More...
CUTLASS_HOST_DEVICE MmaSimtTileIterator & | operator-- ()
| Advances the iterator along the advance dimension. More...
CUTLASS_HOST_DEVICE void | load_with_pointer_offset (Fragment &frag, Index pointer_offset) const
| Loads a fragment from memory at the location pointed to by the iterator. More...
CUTLASS_HOST_DEVICE void | load (Fragment &frag) const
| Loads a fragment from memory at the location pointed to by the iterator. More...
CUTLASS_HOST_DEVICE void | store_with_pointer_offset (Fragment const &frag, Index pointer_offset) const
| Stores a fragment to memory at the location pointed to by the iterator. More...
CUTLASS_HOST_DEVICE void | store (Fragment const &frag) const
| Stores a fragment to memory at the location pointed to by the iterator. More...
CUTLASS_DEVICE void | set_kgroup_index (int k_group)

## Static Public Attributes

static Operand const | kOperand = Operand::kA
| Operand tag. More...
static const int | kInterleave = 4
| Iterleave factor. More...
static const int | kPartitionsK = PartitionsK
| Number of partitions along K dimension. More...
static const int | kGroupPerTile = PartitionGroupSize / Shape::kColumn
| Number of KGroups per kPartition. More...

## Detailed Description

### template<typename Shape_, typename Element_, typename Policy_, int PartitionsK, int PartitionGroupSize>
class cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize >

Specialization for A operands of column-major-K interleaved layouts

Concept: MutableRandomAccessContiguousTileIteratorConcept

## Member Typedef Documentation

template<typename Shape_ , typename Element_ , typename Policy_ , int PartitionsK, int PartitionGroupSize>

using cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize >::Element = Element_

template<typename Shape_ , typename Element_ , typename Policy_ , int PartitionsK, int PartitionGroupSize>

using cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize >::Fragment = Array<Element, ThreadShape::kCount>

template<typename Shape_ , typename Element_ , typename Policy_ , int PartitionsK, int PartitionGroupSize>

using cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize >::Index = typename TensorRef::Index

template<typename Shape_ , typename Element_ , typename Policy_ , int PartitionsK, int PartitionGroupSize>

using cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize >::Iterations = MatrixShape< ThreadShape::kRow / Policy::LaneMmaShape::kM, ThreadShape::kColumn / Policy::LaneMmaShape::kK >

template<typename Shape_ , typename Element_ , typename Policy_ , int PartitionsK, int PartitionGroupSize>

using cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize >::Layout = layout::ColumnMajorInterleaved<4>

template<typename Shape_ , typename Element_ , typename Policy_ , int PartitionsK, int PartitionGroupSize>

using cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize >::LongIndex = typename TensorRef::LongIndex

template<typename Shape_ , typename Element_ , typename Policy_ , int PartitionsK, int PartitionGroupSize>

using cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize >::Policy = Policy_

template<typename Shape_ , typename Element_ , typename Policy_ , int PartitionsK, int PartitionGroupSize>

using cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize >::Shape = Shape_

template<typename Shape_ , typename Element_ , typename Policy_ , int PartitionsK, int PartitionGroupSize>

using cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize >::TensorCoord = typename TensorRef::TensorCoord

template<typename Shape_ , typename Element_ , typename Policy_ , int PartitionsK, int PartitionGroupSize>

using cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize >::TensorRef = TensorRef<Element, Layout>

template<typename Shape_ , typename Element_ , typename Policy_ , int PartitionsK, int PartitionGroupSize>

using cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize >::ThreadShape = MatrixShape< Shape::kRow / Policy::WarpShape::kRow, Shape::kColumn >

## Constructor & Destructor Documentation

template<typename Shape_ , typename Element_ , typename Policy_ , int PartitionsK, int PartitionGroupSize>

CUTLASS_HOST_DEVICE cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize >::MmaSimtTileIterator | ( |  | ) |

inline

template<typename Shape_ , typename Element_ , typename Policy_ , int PartitionsK, int PartitionGroupSize>

CUTLASS_HOST_DEVICE cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize >::MmaSimtTileIterator | ( | TensorRef | ref,
|  | int | lane_id
| ) |  |

inline

## Member Function Documentation

template<typename Shape_ , typename Element_ , typename Policy_ , int PartitionsK, int PartitionGroupSize>

CUTLASS_HOST_DEVICE MmaSimtTileIterator& cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize >::add_pointer_offset | ( | LongIndex | offset | ) |

inline

template<typename Shape_ , typename Element_ , typename Policy_ , int PartitionsK, int PartitionGroupSize>

CUTLASS_HOST_DEVICE MmaSimtTileIterator& cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize >::add_tile_offset | ( | TensorCoord const & | coord | ) |

inline

template<typename Shape_ , typename Element_ , typename Policy_ , int PartitionsK, int PartitionGroupSize>

CUTLASS_HOST_DEVICE void cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize >::load | ( | Fragment & | frag | ) | const

inline

template<typename Shape_ , typename Element_ , typename Policy_ , int PartitionsK, int PartitionGroupSize>

CUTLASS_HOST_DEVICE void cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize >::load_with_pointer_offset | ( | Fragment & | frag,
|  | Index | pointer_offset
| ) |  | const

inline

template<typename Shape_ , typename Element_ , typename Policy_ , int PartitionsK, int PartitionGroupSize>

CUTLASS_HOST_DEVICE MmaSimtTileIterator& cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize >::operator++ | ( |  | ) |

inline

template<typename Shape_ , typename Element_ , typename Policy_ , int PartitionsK, int PartitionGroupSize>

CUTLASS_HOST_DEVICE MmaSimtTileIterator& cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize >::operator-- | ( |  | ) |

inline

template<typename Shape_ , typename Element_ , typename Policy_ , int PartitionsK, int PartitionGroupSize>

CUTLASS_DEVICE void cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize >::set_kgroup_index | ( | int | k_group | ) |

inline

Notify the iterator which k-group it is currently pointing to.

This does not advance the iterator. Rather, it overrides its internal tracking with constant-valued k-group index to enable the compiler to fold constants and achieve more efficient code.

This is used by some nontrivial permuted layouts.

template<typename Shape_ , typename Element_ , typename Policy_ , int PartitionsK, int PartitionGroupSize>

CUTLASS_HOST_DEVICE void cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize >::store | ( | Fragment const & | frag | ) | const

inline

template<typename Shape_ , typename Element_ , typename Policy_ , int PartitionsK, int PartitionGroupSize>

CUTLASS_HOST_DEVICE void cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize >::store_with_pointer_offset | ( | Fragment const & | frag,
|  | Index | pointer_offset
| ) |  | const

inline

## Member Data Documentation

template<typename Shape_ , typename Element_ , typename Policy_ , int PartitionsK, int PartitionGroupSize>

const int cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize >::kGroupPerTile = PartitionGroupSize / Shape::kColumn

static

template<typename Shape_ , typename Element_ , typename Policy_ , int PartitionsK, int PartitionGroupSize>

const int cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize >::kInterleave = 4

static

template<typename Shape_ , typename Element_ , typename Policy_ , int PartitionsK, int PartitionGroupSize>

Operand const cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize >::kOperand = Operand::kA

static

template<typename Shape_ , typename Element_ , typename Policy_ , int PartitionsK, int PartitionGroupSize>

const int cutlass::gemm::warp::MmaSimtTileIterator< Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved< 4 >, Policy_, PartitionsK, PartitionGroupSize >::kPartitionsK = PartitionsK

static

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

- mma_simt_tile_iterator.h