#include <mma_tensor_op_tile_iterator_sm70.h>

Classes
struct	Policy
	Internal structure of iterator - made public to enable introspection. More...

Public Types
using	Shape = Shape_
	Shape of tile to load (concept: MatrixShape) More...

using	Element = Element_
	Element type. More...

using	Layout = Layout_
	Layout of source tile. More...

using	InstructionShape = InstructionShape_
	Shape of one matrix product operation (concept: MatrixShape) More...

using	OpDelta = OpDelta_
	Delta between MMA operations (in units of MMA operations, concept: MatrixShape) 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	Fragment = Array< Element, Shape::kCount/kThreads >
	Fragment object holding a thread's part of a tile. More...

Public Member Functions
CUTLASS_HOST_DEVICE	MmaVoltaTensorOpAccumulatorTileIterator ()
	Default ctor constructs null iterator. More...

CUTLASS_HOST_DEVICE	MmaVoltaTensorOpAccumulatorTileIterator (TensorRef const &ref, int lane_id)
	Constructor from TensorRef. More...

CUTLASS_HOST_DEVICE MmaVoltaTensorOpAccumulatorTileIterator &	add_pointer_offset (LongIndex offset)
	Adds a pointer offset to internal pointer(s) to advance through memory. More...

CUTLASS_HOST_DEVICE MmaVoltaTensorOpAccumulatorTileIterator &	add_tile_offset (TensorCoord const &tile_offset)
	Advances an iterator along logical dimensions of matrix in units of whole tiles. More...

CUTLASS_HOST_DEVICE MmaVoltaTensorOpAccumulatorTileIterator &	operator++ ()
	Advances the iterator along the advance dimension. More...

CUTLASS_HOST_DEVICE MmaVoltaTensorOpAccumulatorTileIterator &	operator-- ()
	Advances the iterator along the advance dimension. More...

CUTLASS_DEVICE MmaVoltaTensorOpAccumulatorTileIterator &	operator+= (TensorCoord const &tile_offset)
	advances in units of whole tiles along the logical coordinate space of the tensor More...

CUTLASS_DEVICE MmaVoltaTensorOpAccumulatorTileIterator &	operator-= (TensorCoord const &tile_offset)

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	load_with_pointer_offset (Fragment &frag, Index pointer_offset) const
	Loads a fragment from memory with additional logical offset. More...

CUTLASS_DEVICE void	load_with_byte_offset (Fragment &frag, Index byte_offset) const
	Loads a fragment from memory with additional logical offset. More...

CUTLASS_HOST_DEVICE void	load (Fragment &frag, TensorCoord const &tile_offset) const
	Loads a fragment from memory with logical offset in units of whole tiles. More...

CUTLASS_HOST_DEVICE void	load (Fragment &frag, TensorCoord const &tile_offset, Index pointer_offset) const
	Loads a fragment from memory with logical offset in units of whole tiles. More...

CUTLASS_HOST_DEVICE void	store (Fragment const &frag) const
	Stores a fragment to memory. More...

CUTLASS_HOST_DEVICE void	store_with_pointer_offset (Fragment const &frag, Index pointer_offset) const
	Stores a fragment to memory with additional pointer offset. More...

CUTLASS_HOST_DEVICE void	store_with_byte_offset (Fragment const &frag, Index byte_offset) const
	Stores a fragment to memory with additional pointer offset. More...

CUTLASS_HOST_DEVICE void	store (Fragment &frag, TensorCoord const &tile_offset) const
	Stores a fragment to memory with logical offset in units of whole tiles. More...

CUTLASS_HOST_DEVICE void	store (Fragment const &frag, TensorCoord const &tile_offset, Index pointer_offset) const
	Stores a fragment from memory with logical offset in units of whole tiles. More...

Static Public Attributes
static Operand const	kOperand = Operand::kC
	Operand tag. More...

static int const	kThreads = 32
	Number of participating threads. More...

Detailed Description

template<typename Shape_, typename Element_, typename Layout_, typename InstructionShape_, typename OpDelta_>
class cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >

This tile iterator is specialized for 32-thread TensorOps. It is used to load or store accumulators from memory and is agnostic to layout. It could be faster if it assumed row-major accumulator layout.

Satisfies: ReadableRandomAccessContiguousTileIteratorConcept | WriteableRandomAccessContiguousTileIteratorConcept

Member Typedef Documentation

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

using cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::Element = Element_

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

using cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::Fragment = Array<Element, Shape::kCount / kThreads>

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

using cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::Index = typename TensorRef::Index

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

using cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::InstructionShape = InstructionShape_

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

using cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::Layout = Layout_

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

using cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::LongIndex = typename TensorRef::LongIndex

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

using cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::OpDelta = OpDelta_

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

using cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::Shape = Shape_

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

using cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::TensorCoord = typename TensorRef::TensorCoord

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

using cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::TensorRef = TensorRef<Element, Layout>

Constructor & Destructor Documentation

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

CUTLASS_HOST_DEVICE cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::MmaVoltaTensorOpAccumulatorTileIterator ( )

inline

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

CUTLASS_HOST_DEVICE cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::MmaVoltaTensorOpAccumulatorTileIterator	(	TensorRef const &	ref,
		int	lane_id
	)

inline

Member Function Documentation

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

CUTLASS_HOST_DEVICE MmaVoltaTensorOpAccumulatorTileIterator& cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::add_pointer_offset ( LongIndex offset )

inline

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

CUTLASS_HOST_DEVICE MmaVoltaTensorOpAccumulatorTileIterator& cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::add_tile_offset ( TensorCoord const & tile_offset )

inline

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

CUTLASS_HOST_DEVICE void cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::load ( Fragment & frag ) const

inline

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

CUTLASS_HOST_DEVICE void cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::load	(	Fragment &	frag,
		TensorCoord const &	tile_offset
	)		const

inline

< loads a tile with a logical offset in units of whole tiles

Parameters

frag	fragment to load from the tensor

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

CUTLASS_HOST_DEVICE void cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::load	(	Fragment &	frag,
		TensorCoord const &	tile_offset,
		Index	pointer_offset
	)		const

inline

< loads a tile with a logical offset AND a pointer offset

Parameters

frag	fragment to load from the tensor
tile_offset	loads a tile with a logical offset in units of whole tiles

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

CUTLASS_DEVICE void cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::load_with_byte_offset	(	Fragment &	frag,
		Index	byte_offset
	)		const

inline

< loads a tile with a linear offset

Parameters

frag	fragment to load from the tensor

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

CUTLASS_HOST_DEVICE void cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::load_with_pointer_offset	(	Fragment &	frag,
		Index	pointer_offset
	)		const

inline

< loads a tile with a linear offset

Parameters

frag	fragment to load from the tensor

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

CUTLASS_HOST_DEVICE MmaVoltaTensorOpAccumulatorTileIterator& cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::operator++ ( )

inline

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

CUTLASS_DEVICE MmaVoltaTensorOpAccumulatorTileIterator& cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::operator+= ( TensorCoord const & tile_offset )

inline

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

CUTLASS_HOST_DEVICE MmaVoltaTensorOpAccumulatorTileIterator& cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::operator-- ( )

inline

advances in units of whole tiles along the logical coordinate space of the tensor

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

CUTLASS_DEVICE MmaVoltaTensorOpAccumulatorTileIterator& cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::operator-= ( TensorCoord const & tile_offset )

inline

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

CUTLASS_HOST_DEVICE void cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::store ( Fragment const & frag ) const

inline

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

CUTLASS_HOST_DEVICE void cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::store	(	Fragment &	frag,
		TensorCoord const &	tile_offset
	)		const

inline

< stores a tile with a logical offset in units of whole tiles

Parameters

frag	fragment to store to the tensor

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

CUTLASS_HOST_DEVICE void cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::store	(	Fragment const &	frag,
		TensorCoord const &	tile_offset,
		Index	pointer_offset
	)		const

inline

Parameters

frag	fragment to store to the tensor
tile_offset	stores a tile with a logical offset in units of whole tiles
pointer_offset	stores a tile with a logical offset AND a pointer offset

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

CUTLASS_HOST_DEVICE void cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::store_with_byte_offset	(	Fragment const &	frag,
		Index	byte_offset
	)		const

inline

< store a tile with a linear offset

Parameters

frag	fragment to store from the tensor

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

CUTLASS_HOST_DEVICE void cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::store_with_pointer_offset	(	Fragment const &	frag,
		Index	pointer_offset
	)		const

inline

< store a tile with a linear offset

Parameters

frag	fragment to store from the tensor

Member Data Documentation

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

Operand const cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::kOperand = Operand::kC

static

template<typename Shape_ , typename Element_ , typename Layout_ , typename InstructionShape_ , typename OpDelta_ >

int const cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >::kThreads = 32

static

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

mma_tensor_op_tile_iterator_sm70.h

Classes

Public Types

Public Member Functions

Static Public Attributes

Detailed Description

template<typename Shape_, typename Element_, typename Layout_, typename InstructionShape_, typename OpDelta_> class cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >

Member Typedef Documentation

Constructor & Destructor Documentation

Member Function Documentation

Member Data Documentation

template<typename Shape_, typename Element_, typename Layout_, typename InstructionShape_, typename OpDelta_>
class cutlass::gemm::warp::MmaVoltaTensorOpAccumulatorTileIterator< Shape_, Element_, Layout_, InstructionShape_, OpDelta_ >