Cutlass
CUDA Templates for Linear Algebra Subroutines and Solvers
Public Types | List of all members
cutlass::gemm::GemmEpilogueTraitsHelper< GemmConfig_, EpilogueFunctor_, Index_ > Struct Template Reference

#include <gemm_epilogue_traits.h>

Public Types

typedef EpilogueFunctor_::Scalar Scalar
 The scalar. More...
 
typedef GemmConfig_::OutputTile OutputTile
 The output tile. More...
 
typedef Shape< 1, GemmConfig_::MultiplyAdd::AccumulatorsPerThread::kH/GemmConfig_::kAccumulatorsPerLdsB, GemmConfig_::kAccumulatorsPerLdsB > Iterations
 The number of iterations in the epilogue. More...
 
typedef Shape< 0, GemmConfig_::kAccumulatorsPerLdsB *(GemmConfig_::Warps::kH *GemmConfig_::MultiplyAdd::ThreadsPerWarp::kH - 1), 0 > Delta
 
typedef EpilogueFunctor_ Functor
 The functor to do the math in the epilogue. More...
 
typedef GemmSharedStoreTileDTraits< typename Functor::ScalarAccum, typename GemmConfig_::OutputTile, typename GemmConfig_::Warps, typename GemmConfig_::MultiplyAdd::ThreadsPerWarp, GemmConfig_::kScalarsPerStsD, 128/sizeof(typename GemmConfig_::ScalarD)/GemmConfig_::kScalarsPerStsD/2 *GemmConfig_::kScalarsPerStsD > SharedStoreTileTraits
 The traits class to build the iterator to store to shared memory for D. More...
 
typedef TileStoreIterator< SharedStoreTileTraits, typename SharedStoreTileTraits::Scalar, IteratorAdvance::kH, MemorySpace::kSharedSharedStoreIteratorD
 The iterator to store D to shared memory. More...
 
typedef Copy< typename SharedStoreIteratorD::FragmentSharedStoreTransformerD
 The shared store transformer for D. More...
 
typedef GemmSharedLoadTileDTraits< typename Functor::ScalarAccum, typename GemmConfig_::OutputTile, typename GemmConfig_::Warps, typename GemmConfig_::MultiplyAdd::ThreadsPerWarp, GemmConfig_::OutputTile::kH/ShapeCount< Iterations >::kCount, GemmConfig_::kScalarsPerLdsD, SharedStoreTileTraits::kSkewSharedLoadTileTraits
 The traits class to build the iterator to load from shared memory for D. More...
 
typedef TileLoadIterator< SharedLoadTileTraits, typename SharedLoadTileTraits::Scalar, IteratorAdvance::kH, MemorySpace::kSharedSharedLoadIteratorD
 The iterator to load D from shared memory. More...
 
typedef SharedLoadStream< SharedLoadIteratorDSharedLoadStreamD
 The stream to load D. More...
 
typedef GemmGlobalTileCdTraits< typename GemmConfig_::ScalarC const, Shape< 1, GemmConfig_::OutputTile::kH/ShapeCount< Iterations >::kCount, GemmConfig_::OutputTile::kW >, Shape< 1, ShapeCount< typename GemmConfig_::Warps >::kCount, GemmConfig_::kWarpSize >, Iterations::kW, GemmConfig_::kScalarsPerLdgC > GlobalLoadTileTraits
 The traits class to build the iterator to load data from global memory for C^N. More...
 
typedef GemmGlobalIteratorCd< GlobalLoadTileTraits, Index_ > GlobalLoadIteratorC
 The iterator to load C. More...
 
typedef Copy< typename GlobalLoadIteratorC::FragmentGlobalTransformerC
 The transformer for C. More...
 
typedef GemmGlobalTileCdTraits< typename GemmConfig_::ScalarD, Shape< 1, GemmConfig_::OutputTile::kH/ShapeCount< Iterations >::kCount, GemmConfig_::OutputTile::kW >, Shape< 1, ShapeCount< typename GemmConfig_::Warps >::kCount, GemmConfig_::kWarpSize >, Iterations::kW, GemmConfig_::kScalarsPerStgD > GlobalStoreTileTraits
 The traits class to build the iterator to store data to global memory for D^N. More...
 
typedef GemmGlobalIteratorCd< GlobalStoreTileTraits, Index_ > GlobalStoreIteratorD
 The iterator to store D. More...
 
typedef Copy< typename GlobalStoreIteratorD::FragmentGlobalTransformerD
 The transformer for D. More...
 

Member Typedef Documentation

◆ Delta

template<typename GemmConfig_, typename EpilogueFunctor_, typename Index_ = int>
typedef Shape<0, GemmConfig_::kAccumulatorsPerLdsB*( GemmConfig_::Warps::kH* GemmConfig_::MultiplyAdd::ThreadsPerWarp::kH - 1), 0> cutlass::gemm::GemmEpilogueTraitsHelper< GemmConfig_, EpilogueFunctor_, Index_ >::Delta

◆ Functor

template<typename GemmConfig_, typename EpilogueFunctor_, typename Index_ = int>
typedef EpilogueFunctor_ cutlass::gemm::GemmEpilogueTraitsHelper< GemmConfig_, EpilogueFunctor_, Index_ >::Functor

◆ GlobalLoadIteratorC

template<typename GemmConfig_, typename EpilogueFunctor_, typename Index_ = int>
typedef GemmGlobalIteratorCd<GlobalLoadTileTraits, Index_> cutlass::gemm::GemmEpilogueTraitsHelper< GemmConfig_, EpilogueFunctor_, Index_ >::GlobalLoadIteratorC

◆ GlobalLoadTileTraits

template<typename GemmConfig_, typename EpilogueFunctor_, typename Index_ = int>
typedef GemmGlobalTileCdTraits< typename GemmConfig_::ScalarC const, Shape<1, GemmConfig_::OutputTile::kH / ShapeCount<Iterations>::kCount, GemmConfig_::OutputTile::kW>, Shape<1, ShapeCount<typename GemmConfig_::Warps>::kCount, GemmConfig_::kWarpSize>, Iterations::kW, GemmConfig_::kScalarsPerLdgC> cutlass::gemm::GemmEpilogueTraitsHelper< GemmConfig_, EpilogueFunctor_, Index_ >::GlobalLoadTileTraits

◆ GlobalStoreIteratorD

template<typename GemmConfig_, typename EpilogueFunctor_, typename Index_ = int>
typedef GemmGlobalIteratorCd<GlobalStoreTileTraits, Index_> cutlass::gemm::GemmEpilogueTraitsHelper< GemmConfig_, EpilogueFunctor_, Index_ >::GlobalStoreIteratorD

◆ GlobalStoreTileTraits

template<typename GemmConfig_, typename EpilogueFunctor_, typename Index_ = int>
typedef GemmGlobalTileCdTraits< typename GemmConfig_::ScalarD, Shape<1, GemmConfig_::OutputTile::kH / ShapeCount<Iterations>::kCount, GemmConfig_::OutputTile::kW>, Shape<1, ShapeCount<typename GemmConfig_::Warps>::kCount, GemmConfig_::kWarpSize>, Iterations::kW, GemmConfig_::kScalarsPerStgD> cutlass::gemm::GemmEpilogueTraitsHelper< GemmConfig_, EpilogueFunctor_, Index_ >::GlobalStoreTileTraits

◆ GlobalTransformerC

template<typename GemmConfig_, typename EpilogueFunctor_, typename Index_ = int>
typedef Copy<typename GlobalLoadIteratorC::Fragment> cutlass::gemm::GemmEpilogueTraitsHelper< GemmConfig_, EpilogueFunctor_, Index_ >::GlobalTransformerC

◆ GlobalTransformerD

template<typename GemmConfig_, typename EpilogueFunctor_, typename Index_ = int>
typedef Copy<typename GlobalStoreIteratorD::Fragment> cutlass::gemm::GemmEpilogueTraitsHelper< GemmConfig_, EpilogueFunctor_, Index_ >::GlobalTransformerD

◆ Iterations

template<typename GemmConfig_, typename EpilogueFunctor_, typename Index_ = int>
typedef Shape<1, GemmConfig_::MultiplyAdd::AccumulatorsPerThread::kH / GemmConfig_::kAccumulatorsPerLdsB, GemmConfig_::kAccumulatorsPerLdsB> cutlass::gemm::GemmEpilogueTraitsHelper< GemmConfig_, EpilogueFunctor_, Index_ >::Iterations

◆ OutputTile

template<typename GemmConfig_, typename EpilogueFunctor_, typename Index_ = int>
typedef GemmConfig_::OutputTile cutlass::gemm::GemmEpilogueTraitsHelper< GemmConfig_, EpilogueFunctor_, Index_ >::OutputTile

◆ Scalar

template<typename GemmConfig_, typename EpilogueFunctor_, typename Index_ = int>
typedef EpilogueFunctor_::Scalar cutlass::gemm::GemmEpilogueTraitsHelper< GemmConfig_, EpilogueFunctor_, Index_ >::Scalar

◆ SharedLoadIteratorD

template<typename GemmConfig_, typename EpilogueFunctor_, typename Index_ = int>
typedef TileLoadIterator<SharedLoadTileTraits, typename SharedLoadTileTraits::Scalar, IteratorAdvance::kH, MemorySpace::kShared> cutlass::gemm::GemmEpilogueTraitsHelper< GemmConfig_, EpilogueFunctor_, Index_ >::SharedLoadIteratorD

◆ SharedLoadStreamD

template<typename GemmConfig_, typename EpilogueFunctor_, typename Index_ = int>
typedef SharedLoadStream<SharedLoadIteratorD> cutlass::gemm::GemmEpilogueTraitsHelper< GemmConfig_, EpilogueFunctor_, Index_ >::SharedLoadStreamD

◆ SharedLoadTileTraits

template<typename GemmConfig_, typename EpilogueFunctor_, typename Index_ = int>
typedef GemmSharedLoadTileDTraits< typename Functor::ScalarAccum, typename GemmConfig_::OutputTile, typename GemmConfig_::Warps, typename GemmConfig_::MultiplyAdd::ThreadsPerWarp, GemmConfig_::OutputTile::kH / ShapeCount<Iterations>::kCount, GemmConfig_::kScalarsPerLdsD, SharedStoreTileTraits::kSkew> cutlass::gemm::GemmEpilogueTraitsHelper< GemmConfig_, EpilogueFunctor_, Index_ >::SharedLoadTileTraits

◆ SharedStoreIteratorD

template<typename GemmConfig_, typename EpilogueFunctor_, typename Index_ = int>
typedef TileStoreIterator<SharedStoreTileTraits, typename SharedStoreTileTraits::Scalar, IteratorAdvance::kH, MemorySpace::kShared> cutlass::gemm::GemmEpilogueTraitsHelper< GemmConfig_, EpilogueFunctor_, Index_ >::SharedStoreIteratorD

◆ SharedStoreTileTraits

template<typename GemmConfig_, typename EpilogueFunctor_, typename Index_ = int>
typedef GemmSharedStoreTileDTraits< typename Functor::ScalarAccum, typename GemmConfig_::OutputTile, typename GemmConfig_::Warps, typename GemmConfig_::MultiplyAdd::ThreadsPerWarp, GemmConfig_::kScalarsPerStsD, 128 / sizeof(typename GemmConfig_::ScalarD) / GemmConfig_::kScalarsPerStsD / 2 * GemmConfig_::kScalarsPerStsD> cutlass::gemm::GemmEpilogueTraitsHelper< GemmConfig_, EpilogueFunctor_, Index_ >::SharedStoreTileTraits

◆ SharedStoreTransformerD

template<typename GemmConfig_, typename EpilogueFunctor_, typename Index_ = int>
typedef Copy<typename SharedStoreIteratorD::Fragment> cutlass::gemm::GemmEpilogueTraitsHelper< GemmConfig_, EpilogueFunctor_, Index_ >::SharedStoreTransformerD

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