cutlass::epilogue::threadblock::EpilogueBase< Shape_, WarpMmaOperator_, PartitionsK, AccumulatorFragmentIterator_, WarpTileIterator_, Padding_ >::SharedStorage Struct Reference

Shared storage allocation needed by the epilogue.

#include <epilogue_base.h>

Collaboration diagram for cutlass::epilogue::threadblock::EpilogueBase< Shape_, WarpMmaOperator_, PartitionsK, AccumulatorFragmentIterator_, WarpTileIterator_, Padding_ >::SharedStorage:

Public Types
using	Element = typename WarpTileIterator::Element
	Element type of shared memory. More...
using	TensorRef = typename WarpTileIterator::TensorRef
	Tensor reference to shared memory allocation. More...
using	Layout = typename WarpTileIterator::Layout
	Layout of shared memory allocation. More...
using	Shape = MatrixShape< WarpCount::kM *WarpTileIterator::Shape::kRow *WarpCount::kK, WarpCount::kN *WarpTileIterator::Shape::kColumn >
	Logical shape of the shared memory tile written to by all warps. More...
using	StorageShape = MatrixShape< Shape::kRow+Padding::kRow, Shape::kColumn+Padding::kColumn >
	Shape of the shared memory allocation for the epilogue. More...

Public Member Functions
CUTLASS_DEVICE Element *	data ()
	Returns a pointer to the shared memory buffer. More...
CUTLASS_DEVICE TensorRef	reference ()
	Returns a tensor reference to the shared memory buffer. More...
CUTLASS_DEVICE void	debug_print ()

Public Attributes
AlignedBuffer< Element, StorageShape::kCount >	storage

Member Typedef Documentation

template<typename Shape_ , typename WarpMmaOperator_ , int PartitionsK, typename AccumulatorFragmentIterator_ , typename WarpTileIterator_ , typename Padding_ >

using cutlass::epilogue::threadblock::EpilogueBase< Shape_, WarpMmaOperator_, PartitionsK, AccumulatorFragmentIterator_, WarpTileIterator_, Padding_ >::SharedStorage::Element = typename WarpTileIterator::Element

template<typename Shape_ , typename WarpMmaOperator_ , int PartitionsK, typename AccumulatorFragmentIterator_ , typename WarpTileIterator_ , typename Padding_ >

using cutlass::epilogue::threadblock::EpilogueBase< Shape_, WarpMmaOperator_, PartitionsK, AccumulatorFragmentIterator_, WarpTileIterator_, Padding_ >::SharedStorage::Layout = typename WarpTileIterator::Layout

template<typename Shape_ , typename WarpMmaOperator_ , int PartitionsK, typename AccumulatorFragmentIterator_ , typename WarpTileIterator_ , typename Padding_ >

using cutlass::epilogue::threadblock::EpilogueBase< Shape_, WarpMmaOperator_, PartitionsK, AccumulatorFragmentIterator_, WarpTileIterator_, Padding_ >::SharedStorage::Shape = MatrixShape< WarpCount::kM * WarpTileIterator::Shape::kRow * WarpCount::kK, WarpCount::kN * WarpTileIterator::Shape::kColumn >

template<typename Shape_ , typename WarpMmaOperator_ , int PartitionsK, typename AccumulatorFragmentIterator_ , typename WarpTileIterator_ , typename Padding_ >

using cutlass::epilogue::threadblock::EpilogueBase< Shape_, WarpMmaOperator_, PartitionsK, AccumulatorFragmentIterator_, WarpTileIterator_, Padding_ >::SharedStorage::StorageShape = MatrixShape< Shape::kRow + Padding::kRow, Shape::kColumn + Padding::kColumn >

template<typename Shape_ , typename WarpMmaOperator_ , int PartitionsK, typename AccumulatorFragmentIterator_ , typename WarpTileIterator_ , typename Padding_ >

using cutlass::epilogue::threadblock::EpilogueBase< Shape_, WarpMmaOperator_, PartitionsK, AccumulatorFragmentIterator_, WarpTileIterator_, Padding_ >::SharedStorage::TensorRef = typename WarpTileIterator::TensorRef

Member Function Documentation

template<typename Shape_ , typename WarpMmaOperator_ , int PartitionsK, typename AccumulatorFragmentIterator_ , typename WarpTileIterator_ , typename Padding_ >

CUTLASS_DEVICE Element* cutlass::epilogue::threadblock::EpilogueBase< Shape_, WarpMmaOperator_, PartitionsK, AccumulatorFragmentIterator_, WarpTileIterator_, Padding_ >::SharedStorage::data ( )

inline

template<typename Shape_ , typename WarpMmaOperator_ , int PartitionsK, typename AccumulatorFragmentIterator_ , typename WarpTileIterator_ , typename Padding_ >

CUTLASS_DEVICE void cutlass::epilogue::threadblock::EpilogueBase< Shape_, WarpMmaOperator_, PartitionsK, AccumulatorFragmentIterator_, WarpTileIterator_, Padding_ >::SharedStorage::debug_print ( )

inline

template<typename Shape_ , typename WarpMmaOperator_ , int PartitionsK, typename AccumulatorFragmentIterator_ , typename WarpTileIterator_ , typename Padding_ >

CUTLASS_DEVICE TensorRef cutlass::epilogue::threadblock::EpilogueBase< Shape_, WarpMmaOperator_, PartitionsK, AccumulatorFragmentIterator_, WarpTileIterator_, Padding_ >::SharedStorage::reference ( )

inline

Member Data Documentation

template<typename Shape_ , typename WarpMmaOperator_ , int PartitionsK, typename AccumulatorFragmentIterator_ , typename WarpTileIterator_ , typename Padding_ >

AlignedBuffer<Element, StorageShape::kCount> cutlass::epilogue::threadblock::EpilogueBase< Shape_, WarpMmaOperator_, PartitionsK, AccumulatorFragmentIterator_, WarpTileIterator_, Padding_ >::SharedStorage::storage

---

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

• epilogue_base.h