CUTLASS: cutlass::epilogue::threadblock::EpilogueBase< Shape_, WarpMmaOperator_, PartitionsK, AccumulatorFragmentIterator_, WarpTileIterator_, Padding_ > Class Template Reference

CUTLASS

CUDA Templates for Linear Algebra Subroutines and Solvers

Base class for epilogues defining warp-level.

#include <epilogue_base.h>

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

Inheritance graph

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

Collaboration graph

Classes
struct	SharedStorage
	Shared storage allocation needed by the epilogue. More...

Public Types
using	Shape = Shape_

using	WarpMmaOperator = WarpMmaOperator_

using	AccumulatorFragmentIterator = AccumulatorFragmentIterator_

using	WarpTileIterator = WarpTileIterator_

using	Padding = Padding_

using	Layout = layout::RowMajor
	Output layout is always row-major. More...

using	AccumulatorTile = typename AccumulatorFragmentIterator::AccumulatorTile
	The complete warp-level accumulator tile. More...

using	ElementAccumulator = typename AccumulatorTile::Element
	Accumulator element. More...

using	WarpCount = gemm::GemmShape< Shape::kM/WarpMmaOperator::Shape::kM, Shape::kN/WarpMmaOperator::Shape::kN, kPartitionsK >
	Number of warps. More...

Public Member Functions
CUTLASS_DEVICE	EpilogueBase (SharedStorage &shared_storage, int thread_idx, int warp_idx, int lane_idx)
	Constructor. More...

Static Public Attributes
static int const	kPartitionsK = PartitionsK

Protected Attributes
SharedStorage &	shared_storage_

WarpTileIterator	warp_tile_iterator_
	Stores a warp's fragment of accumulators to SMEM. More...

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_ >::AccumulatorFragmentIterator = AccumulatorFragmentIterator_

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

using cutlass::epilogue::threadblock::EpilogueBase< Shape_, WarpMmaOperator_, PartitionsK, AccumulatorFragmentIterator_, WarpTileIterator_, Padding_ >::AccumulatorTile = typename AccumulatorFragmentIterator::AccumulatorTile

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

using cutlass::epilogue::threadblock::EpilogueBase< Shape_, WarpMmaOperator_, PartitionsK, AccumulatorFragmentIterator_, WarpTileIterator_, Padding_ >::ElementAccumulator = typename AccumulatorTile::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_ >::Layout = layout::RowMajor

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

using cutlass::epilogue::threadblock::EpilogueBase< Shape_, WarpMmaOperator_, PartitionsK, AccumulatorFragmentIterator_, WarpTileIterator_, Padding_ >::Padding = Padding_

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

using cutlass::epilogue::threadblock::EpilogueBase< Shape_, WarpMmaOperator_, PartitionsK, AccumulatorFragmentIterator_, WarpTileIterator_, Padding_ >::Shape = Shape_

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

using cutlass::epilogue::threadblock::EpilogueBase< Shape_, WarpMmaOperator_, PartitionsK, AccumulatorFragmentIterator_, WarpTileIterator_, Padding_ >::WarpCount = gemm::GemmShape< Shape::kM / WarpMmaOperator::Shape::kM, Shape::kN / WarpMmaOperator::Shape::kN, kPartitionsK >

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

using cutlass::epilogue::threadblock::EpilogueBase< Shape_, WarpMmaOperator_, PartitionsK, AccumulatorFragmentIterator_, WarpTileIterator_, Padding_ >::WarpMmaOperator = WarpMmaOperator_

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

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

Constructor & Destructor Documentation

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

CUTLASS_DEVICE cutlass::epilogue::threadblock::EpilogueBase< Shape_, WarpMmaOperator_, PartitionsK, AccumulatorFragmentIterator_, WarpTileIterator_, Padding_ >::EpilogueBase	(	SharedStorage &	shared_storage,
		int	thread_idx,
		int	warp_idx,
		int	lane_idx
	)

inline

Parameters

shared_storage	Shared storage object
thread_idx	ID of a thread within the threadblock
warp_idx	ID of warp within threadblock
lane_idx	Id of thread within warp

Member Data Documentation

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

int const cutlass::epilogue::threadblock::EpilogueBase< Shape_, WarpMmaOperator_, PartitionsK, AccumulatorFragmentIterator_, WarpTileIterator_, Padding_ >::kPartitionsK = PartitionsK

static

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

SharedStorage& cutlass::epilogue::threadblock::EpilogueBase< Shape_, WarpMmaOperator_, PartitionsK, AccumulatorFragmentIterator_, WarpTileIterator_, Padding_ >::shared_storage_

protected

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

WarpTileIterator cutlass::epilogue::threadblock::EpilogueBase< Shape_, WarpMmaOperator_, PartitionsK, AccumulatorFragmentIterator_, WarpTileIterator_, Padding_ >::warp_tile_iterator_

protected

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

epilogue_base.h