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

Base class for epilogues defining warp-level.

`#include <epilogue_base.h>`

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

## 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