cutlass::epilogue::threadblock::DirectEpilogueTensorOp< Shape_, Operator_, PartitionsK, Element_, OutputOp_, ConvertOp_ > Class Template Reference

Epilogue operator.

#include <direct_epilogue_tensor_op.h>

Classes
struct	Params
	Parameters structure for host-constructible state. More...
struct	SharedStorage
	Shared storage allocation needed by the epilogue. More...

Public Types
using	Shape = Shape_
using	Operator = Operator_
using	WarpCount = gemm::GemmShape< Shape::kM/Operator::Shape::kM, Shape::kN/Operator::Shape::kN, PartitionsK, >
	Number of warps spanning threadblock-scoped tile. More...
using	FragmentC = typename Operator::FragmentC
	Accumulator tile is really the warp-scoped tile. More...
using	Element = Element_
	Data type of output tensor. More...
using	Layout = layout::RowMajor
	Output layout is always row-major. More...
using	OutputOp = OutputOp_
	Function operator computing final output. More...
using	ConvertOp = ConvertOp_
	Conversion operator to shared memory. More...
using	TensorRef = TensorRef< Element, Layout::kRank, Layout >
	Reference to source and destination tensors. More...

Public Member Functions
CUTLASS_DEVICE	DirectEpilogueTensorOp (Params const &params, SharedStorage &shared_storage, int thread_idx, int warp_idx, int lane_idx)
	Constructor. More...
CUTLASS_DEVICE void	operator() (gemm::GemmCoord problem_size, gemm::GemmCoord tb_tile_coord, FragmentC const &accumulators)
	Streams the result to global memory. More...

Member Typedef Documentation

template<typename Shape_ , typename Operator_ , int PartitionsK, typename Element_ , typename OutputOp_ , typename ConvertOp_ >

using cutlass::epilogue::threadblock::DirectEpilogueTensorOp< Shape_, Operator_, PartitionsK, Element_, OutputOp_, ConvertOp_ >::ConvertOp = ConvertOp_

template<typename Shape_ , typename Operator_ , int PartitionsK, typename Element_ , typename OutputOp_ , typename ConvertOp_ >

using cutlass::epilogue::threadblock::DirectEpilogueTensorOp< Shape_, Operator_, PartitionsK, Element_, OutputOp_, ConvertOp_ >::Element = Element_

template<typename Shape_ , typename Operator_ , int PartitionsK, typename Element_ , typename OutputOp_ , typename ConvertOp_ >

using cutlass::epilogue::threadblock::DirectEpilogueTensorOp< Shape_, Operator_, PartitionsK, Element_, OutputOp_, ConvertOp_ >::FragmentC = typename Operator::FragmentC

template<typename Shape_ , typename Operator_ , int PartitionsK, typename Element_ , typename OutputOp_ , typename ConvertOp_ >

using cutlass::epilogue::threadblock::DirectEpilogueTensorOp< Shape_, Operator_, PartitionsK, Element_, OutputOp_, ConvertOp_ >::Layout = layout::RowMajor

template<typename Shape_ , typename Operator_ , int PartitionsK, typename Element_ , typename OutputOp_ , typename ConvertOp_ >

using cutlass::epilogue::threadblock::DirectEpilogueTensorOp< Shape_, Operator_, PartitionsK, Element_, OutputOp_, ConvertOp_ >::Operator = Operator_

template<typename Shape_ , typename Operator_ , int PartitionsK, typename Element_ , typename OutputOp_ , typename ConvertOp_ >

using cutlass::epilogue::threadblock::DirectEpilogueTensorOp< Shape_, Operator_, PartitionsK, Element_, OutputOp_, ConvertOp_ >::OutputOp = OutputOp_

template<typename Shape_ , typename Operator_ , int PartitionsK, typename Element_ , typename OutputOp_ , typename ConvertOp_ >

using cutlass::epilogue::threadblock::DirectEpilogueTensorOp< Shape_, Operator_, PartitionsK, Element_, OutputOp_, ConvertOp_ >::Shape = Shape_

template<typename Shape_ , typename Operator_ , int PartitionsK, typename Element_ , typename OutputOp_ , typename ConvertOp_ >

using cutlass::epilogue::threadblock::DirectEpilogueTensorOp< Shape_, Operator_, PartitionsK, Element_, OutputOp_, ConvertOp_ >::TensorRef = TensorRef<Element, Layout::kRank, Layout>

template<typename Shape_ , typename Operator_ , int PartitionsK, typename Element_ , typename OutputOp_ , typename ConvertOp_ >

using cutlass::epilogue::threadblock::DirectEpilogueTensorOp< Shape_, Operator_, PartitionsK, Element_, OutputOp_, ConvertOp_ >::WarpCount = gemm::GemmShape< Shape::kM / Operator::Shape::kM, Shape::kN / Operator::Shape::kN, PartitionsK, >

Constructor & Destructor Documentation

template<typename Shape_ , typename Operator_ , int PartitionsK, typename Element_ , typename OutputOp_ , typename ConvertOp_ >

CUTLASS_DEVICE cutlass::epilogue::threadblock::DirectEpilogueTensorOp< Shape_, Operator_, PartitionsK, Element_, OutputOp_, ConvertOp_ >::DirectEpilogueTensorOp ( Params const &  params, SharedStorage &  shared_storage, int  thread_idx, int  warp_idx, int  lane_idx  )

inline

Parameters

params	Host-constructable params object
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 Function Documentation

template<typename Shape_ , typename Operator_ , int PartitionsK, typename Element_ , typename OutputOp_ , typename ConvertOp_ >

CUTLASS_DEVICE void cutlass::epilogue::threadblock::DirectEpilogueTensorOp< Shape_, Operator_, PartitionsK, Element_, OutputOp_, ConvertOp_ >::operator() ( gemm::GemmCoord  problem_size, gemm::GemmCoord  tb_tile_coord, FragmentC const &  accumulators  )

inline

< Accumulator tile

Number of mma operations performed

Parameters

problem_size	Problem size of GEMM (units of ElementC)
tb_tile_coord	Threadblock tile coordinate in GEMM (in units of threadblock tiles)

---

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

• direct_epilogue_tensor_op.h