CUTLASS
CUDA Templates for Linear Algebra Subroutines and Solvers
Classes | Public Types | Public Member Functions | Static Public Attributes | List of all members
cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 > Class Template Reference

#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: PitchLinearShape) More...
 
using Element = Element_
 Element type. More...
 
using Layout = cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >
 Layout of source tile. More...
 
using InstructionShape = InstructionShape_
 Shape of one matrix product operation (concept: GemmShape) 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 *2 >
 Fragment object holding a thread's part of a tile. More...
 

Public Member Functions

CUTLASS_HOST_DEVICE MmaVoltaTensorOpMultiplicandTileIterator ()
 Default ctor constructs null iterator. More...
 
CUTLASS_DEVICE MmaVoltaTensorOpMultiplicandTileIterator (TensorRef const &ref, int lane_id)
 Constructor from TensorRef. More...
 
CUTLASS_DEVICE MmaVoltaTensorOpMultiplicandTileIteratoradd_pointer_offset (LongIndex offset)
 Adds a pointer offset to internal pointer(s) to advance through memory. More...
 
CUTLASS_HOST_DEVICE MmaVoltaTensorOpMultiplicandTileIteratoradd_tile_offset (TensorCoord const &tile_offset)
 Advances an iterator along logical dimensions of matrix in units of whole tiles. More...
 
CUTLASS_DEVICE MmaVoltaTensorOpMultiplicandTileIteratoroperator++ ()
 Advances the iterator along the advance dimension. More...
 
CUTLASS_HOST_DEVICE MmaVoltaTensorOpMultiplicandTileIteratoroperator-- ()
 Advances the iterator along the advance dimension. More...
 
CUTLASS_DEVICE MmaVoltaTensorOpMultiplicandTileIteratoroperator+= (TensorCoord const &tile_offset)
 advances in units of whole tiles along the logical coordinate space of the tensor More...
 
CUTLASS_DEVICE MmaVoltaTensorOpMultiplicandTileIteratoroperator-= (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_DEVICE void load_with_byte_offset (Fragment &frag, Index byte_offset) const
 Loads a fragment from memory with additional logical offset. More...
 
CUTLASS_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 (Fragment &frag, TensorCoord const &tile_offset) const
 Loads a fragment from memory with logical offset in units of whole tiles. More...
 
CUTLASS_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_DEVICE void load_with_byte_offset (Fragment &frag, TensorCoord const &tile_offset, Index byte_offset) const
 Loads a fragment from memory with logical offset in units of whole tiles. More...
 
CUTLASS_DEVICE void set_kgroup_index (int k_group)
 

Static Public Attributes

static Operand const kOperand = Operand::kA
 Operand tag. More...
 
static int const kOpDelta = OpDelta_
 Delta between *MMA operations (in units of *MMA operations, concept: MatrixShape) More...
 
static int const kThreads = 32
 Number of participating threads. More...
 

Detailed Description

template<typename Shape_, typename Element_, typename InstructionShape_, int OpDelta_>
class cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >

This tile iterator is specialized for 32-thread TensorOps.

Satisfies: ReadableRandomAccessContiguousTileIteratorConcept

Member Typedef Documentation

template<typename Shape_ , typename Element_ , typename InstructionShape_ , int OpDelta_>
using cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >::Element = Element_
template<typename Shape_ , typename Element_ , typename InstructionShape_ , int OpDelta_>
using cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >::Fragment = Array<Element, Shape::kCount / kThreads * 2>
template<typename Shape_ , typename Element_ , typename InstructionShape_ , int OpDelta_>
using cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >::Index = typename TensorRef::Index
template<typename Shape_ , typename Element_ , typename InstructionShape_ , int OpDelta_>
using cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >::InstructionShape = InstructionShape_
template<typename Shape_ , typename Element_ , typename InstructionShape_ , int OpDelta_>
using cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >::Layout = cutlass::layout::VoltaTensorOpMultiplicandCongruous<sizeof_bits<Element_>::value>
template<typename Shape_ , typename Element_ , typename InstructionShape_ , int OpDelta_>
using cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >::LongIndex = typename TensorRef::LongIndex
template<typename Shape_ , typename Element_ , typename InstructionShape_ , int OpDelta_>
using cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >::Shape = Shape_
template<typename Shape_ , typename Element_ , typename InstructionShape_ , int OpDelta_>
using cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >::TensorCoord = typename TensorRef::TensorCoord
template<typename Shape_ , typename Element_ , typename InstructionShape_ , int OpDelta_>
using cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >::TensorRef = TensorRef<Element, Layout>

Constructor & Destructor Documentation

template<typename Shape_ , typename Element_ , typename InstructionShape_ , int OpDelta_>
CUTLASS_HOST_DEVICE cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >::MmaVoltaTensorOpMultiplicandTileIterator ( )
inline
template<typename Shape_ , typename Element_ , typename InstructionShape_ , int OpDelta_>
CUTLASS_DEVICE cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >::MmaVoltaTensorOpMultiplicandTileIterator ( TensorRef const &  ref,
int  lane_id 
)
inline

Member Function Documentation

template<typename Shape_ , typename Element_ , typename InstructionShape_ , int OpDelta_>
CUTLASS_DEVICE MmaVoltaTensorOpMultiplicandTileIterator& cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >::add_pointer_offset ( LongIndex  offset)
inline
template<typename Shape_ , typename Element_ , typename InstructionShape_ , int OpDelta_>
CUTLASS_HOST_DEVICE MmaVoltaTensorOpMultiplicandTileIterator& cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >::add_tile_offset ( TensorCoord const &  tile_offset)
inline
template<typename Shape_ , typename Element_ , typename InstructionShape_ , int OpDelta_>
CUTLASS_HOST_DEVICE void cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >::load ( Fragment frag) const
inline
template<typename Shape_ , typename Element_ , typename InstructionShape_ , int OpDelta_>
CUTLASS_DEVICE void cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >::load ( Fragment frag,
TensorCoord const &  tile_offset 
) const
inline
Parameters
fragfragment to load from the tensor
tile_offsetloads a tile with a logical offset in units of whole tiles
template<typename Shape_ , typename Element_ , typename InstructionShape_ , int OpDelta_>
CUTLASS_DEVICE void cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >::load ( Fragment frag,
TensorCoord const &  tile_offset,
Index  pointer_offset 
) const
inline
Parameters
fragfragment to load from the tensor
tile_offsetloads a tile with a logical offset in units of whole tiles
pointer_offsetloads a tile with a logical offset AND a pointer offset
template<typename Shape_ , typename Element_ , typename InstructionShape_ , int OpDelta_>
CUTLASS_DEVICE void cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >::load_with_byte_offset ( Fragment frag,
Index  byte_offset 
) const
inline
Parameters
fragfragment to load from the tensor
byte_offsetloads a tile with a linear offset in units of bytes
template<typename Shape_ , typename Element_ , typename InstructionShape_ , int OpDelta_>
CUTLASS_DEVICE void cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >::load_with_byte_offset ( Fragment frag,
TensorCoord const &  tile_offset,
Index  byte_offset 
) const
inline
Parameters
fragfragment to load from the tensor
tile_offsetloads a tile with a logical offset in units of whole tiles
byte_offsetloads a tile with a logical offset AND a pointer offset
template<typename Shape_ , typename Element_ , typename InstructionShape_ , int OpDelta_>
CUTLASS_DEVICE void cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >::load_with_pointer_offset ( Fragment frag,
Index  pointer_offset 
) const
inline
Parameters
fragfragment to load from the tensor
pointer_offsetloads a tile with a linear offset
template<typename Shape_ , typename Element_ , typename InstructionShape_ , int OpDelta_>
CUTLASS_DEVICE MmaVoltaTensorOpMultiplicandTileIterator& cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >::operator++ ( )
inline
template<typename Shape_ , typename Element_ , typename InstructionShape_ , int OpDelta_>
CUTLASS_DEVICE MmaVoltaTensorOpMultiplicandTileIterator& cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >::operator+= ( TensorCoord const &  tile_offset)
inline
template<typename Shape_ , typename Element_ , typename InstructionShape_ , int OpDelta_>
CUTLASS_HOST_DEVICE MmaVoltaTensorOpMultiplicandTileIterator& cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >::operator-- ( )
inline

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

template<typename Shape_ , typename Element_ , typename InstructionShape_ , int OpDelta_>
CUTLASS_DEVICE MmaVoltaTensorOpMultiplicandTileIterator& cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >::operator-= ( TensorCoord const &  tile_offset)
inline
template<typename Shape_ , typename Element_ , typename InstructionShape_ , int OpDelta_>
CUTLASS_DEVICE void cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >::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.

Member Data Documentation

template<typename Shape_ , typename Element_ , typename InstructionShape_ , int OpDelta_>
int const cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >::kOpDelta = OpDelta_
static
template<typename Shape_ , typename Element_ , typename InstructionShape_ , int OpDelta_>
Operand const cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >::kOperand = Operand::kA
static
template<typename Shape_ , typename Element_ , typename InstructionShape_ , int OpDelta_>
int const cutlass::gemm::warp::MmaVoltaTensorOpMultiplicandTileIterator< Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous< sizeof_bits< Element_ >::value >, InstructionShape_, OpDelta_, 32 >::kThreads = 32
static

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