CUTLASS
CUDA Templates for Linear Algebra Subroutines and Solvers
Classes | Public Types | Public Member Functions | Static Public Attributes | List of all members
cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor > Class Template Reference

Template for reading and writing tiles of accumulators to shared memory.

#include <tile_iterator_tensor_op.h>

Classes

struct  Detail
 

Public Types

using WarpShape = WarpShape_
 
using OperatorShape = OperatorShape_
 
using Element = Element_
 
using Layout = layout::RowMajor
 
using TensorRef = TensorRef< Element, Layout >
 Tensor Reference object. More...
 
using TensorCoord = MatrixCoord
 Logical coordinate in referenced tensor. More...
 
using Index = typename TensorRef::Index
 
using LongIndex = typename TensorRef::LongIndex
 
using Policy = TensorOpPolicy< WarpShape, OperatorShape, Layout >
 
using Shape = MatrixShape< Policy::kRowsPerIteration, WarpShape::kN >
 Shape of the tile in memory. More...
 
using Fragment = Array< Element, Policy::OperatorCount::kColumn *Policy::kElementsPerAccess >
 This is the fragment size produced by one access of the iterator. More...
 
using Padding = MatrixShape< 0, Detail::kLanesInQuad *Policy::kElementsPerAccess >
 Padding quantity. More...
 

Public Member Functions

CUTLASS_HOST_DEVICE TileIteratorTensorOp ()
 Default constructor. More...
 
CUTLASS_HOST_DEVICE TileIteratorTensorOp (TensorRef const &ref, unsigned lane_id)
 Constructor from TensorRef. More...
 
CUTLASS_HOST_DEVICE TileIteratorTensorOpadd_pointer_offset (Index pointer_offset)
 Adds a pointer offset. More...
 
CUTLASS_HOST_DEVICE TileIteratorTensorOpadd_tile_offset (TensorCoord const &tile_offset)
 advances in units of whole tiles along the logical coordinate space of the tensor More...
 
CUTLASS_HOST_DEVICE TileIteratorTensorOpoperator+= (TensorCoord const &tile_offset)
 
CUTLASS_HOST_DEVICE void store_with_pointer_offset (Fragment const &frag, Index pointer_offset)
 Store. More...
 
CUTLASS_HOST_DEVICE void store (Fragment const &frag)
 Store. More...
 
CUTLASS_HOST_DEVICE void load_with_pointer_offset (Fragment &frag, Index pointer_offset) const
 Load. More...
 
CUTLASS_HOST_DEVICE void load (Fragment &frag) const
 Load. More...
 

Static Public Attributes

static int const kIterations = Policy::kIterations
 This is the complete warp-level accumulator tile. More...
 

Member Typedef Documentation

template<typename WarpShape_ , typename OperatorShape_ , typename Element_ >
using cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::Element = Element_
template<typename WarpShape_ , typename OperatorShape_ , typename Element_ >
using cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::Fragment = Array< Element, Policy::OperatorCount::kColumn * Policy::kElementsPerAccess>
template<typename WarpShape_ , typename OperatorShape_ , typename Element_ >
using cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::Index = typename TensorRef::Index
template<typename WarpShape_ , typename OperatorShape_ , typename Element_ >
using cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::Layout = layout::RowMajor
template<typename WarpShape_ , typename OperatorShape_ , typename Element_ >
using cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::LongIndex = typename TensorRef::LongIndex
template<typename WarpShape_ , typename OperatorShape_ , typename Element_ >
using cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::OperatorShape = OperatorShape_
template<typename WarpShape_ , typename OperatorShape_ , typename Element_ >
using cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::Padding = MatrixShape< 0, Detail::kLanesInQuad * Policy::kElementsPerAccess>
template<typename WarpShape_ , typename OperatorShape_ , typename Element_ >
using cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::Policy = TensorOpPolicy<WarpShape, OperatorShape, Layout>
template<typename WarpShape_ , typename OperatorShape_ , typename Element_ >
using cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::Shape = MatrixShape< Policy::kRowsPerIteration, WarpShape::kN >
template<typename WarpShape_ , typename OperatorShape_ , typename Element_ >
using cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::TensorCoord = MatrixCoord
template<typename WarpShape_ , typename OperatorShape_ , typename Element_ >
using cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::TensorRef = TensorRef<Element, Layout>
template<typename WarpShape_ , typename OperatorShape_ , typename Element_ >
using cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::WarpShape = WarpShape_

Constructor & Destructor Documentation

template<typename WarpShape_ , typename OperatorShape_ , typename Element_ >
CUTLASS_HOST_DEVICE cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::TileIteratorTensorOp ( )
inline
template<typename WarpShape_ , typename OperatorShape_ , typename Element_ >
CUTLASS_HOST_DEVICE cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::TileIteratorTensorOp ( TensorRef const &  ref,
unsigned  lane_id 
)
inline

Member Function Documentation

template<typename WarpShape_ , typename OperatorShape_ , typename Element_ >
CUTLASS_HOST_DEVICE TileIteratorTensorOp& cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::add_pointer_offset ( Index  pointer_offset)
inline

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

template<typename WarpShape_ , typename OperatorShape_ , typename Element_ >
CUTLASS_HOST_DEVICE TileIteratorTensorOp& cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::add_tile_offset ( TensorCoord const &  tile_offset)
inline
template<typename WarpShape_ , typename OperatorShape_ , typename Element_ >
CUTLASS_HOST_DEVICE void cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::load ( Fragment frag) const
inline
template<typename WarpShape_ , typename OperatorShape_ , typename Element_ >
CUTLASS_HOST_DEVICE void cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::load_with_pointer_offset ( Fragment frag,
Index  pointer_offset 
) const
inline
template<typename WarpShape_ , typename OperatorShape_ , typename Element_ >
CUTLASS_HOST_DEVICE TileIteratorTensorOp& cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::operator+= ( TensorCoord const &  tile_offset)
inline
template<typename WarpShape_ , typename OperatorShape_ , typename Element_ >
CUTLASS_HOST_DEVICE void cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::store ( Fragment const &  frag)
inline
template<typename WarpShape_ , typename OperatorShape_ , typename Element_ >
CUTLASS_HOST_DEVICE void cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::store_with_pointer_offset ( Fragment const &  frag,
Index  pointer_offset 
)
inline

Member Data Documentation

template<typename WarpShape_ , typename OperatorShape_ , typename Element_ >
int const cutlass::epilogue::warp::TileIteratorTensorOp< WarpShape_, OperatorShape_, Element_, layout::RowMajor >::kIterations = Policy::kIterations
static

Number of times this iterator can be incremented


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