Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ > Struct Template Reference

Layernorm2dFwdPipelineTwoPass&lt; Problem_, Policy_ &gt; Struct Template Reference#

Composable Kernel: ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ > Struct Template Reference
Public Types | Public Member Functions | Static Public Member Functions | Static Public Attributes | List of all members
ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ > Struct Template Reference

#include <layernorm2d_fwd_pipeline_two_pass.hpp>

Public Types

using Problem = ck_tile::remove_cvref_t<Problem_>
using Policy = ck_tile::remove_cvref_t<Policy_>
using XDataType = ck_tile::remove_cvref_t<typename Problem::XDataType>
using XBiasDataType = ck_tile::remove_cvref_t<typename Problem::XBiasDataType>
using GammaDataType = ck_tile::remove_cvref_t<typename Problem::GammaDataType>
using BetaDataType = ck_tile::remove_cvref_t<typename Problem::BetaDataType>
using ComputeDataType = ck_tile::remove_cvref_t<typename Problem::ComputeDataType>
using YDataType = ck_tile::remove_cvref_t<typename Problem::YDataType>
using MeanDataType = ck_tile::remove_cvref_t<typename Problem::MeanDataType>
using InvStdDataType = ck_tile::remove_cvref_t<typename Problem::InvStdDataType>
using XResidualDataType = XDataType
using YResidualDataType = XDataType

Public Member Functions

template<typename XWindow, typename XResidualWindow, typename XBiasWindow, typename GammaWindow, typename BetaWindow, typename YWindow, typename YResidualWindow, typename MeanWindow, typename InvStdWindow, typename SmoothScaleWindow, typename YScaleWindow, typename Epilogue>
CK_TILE_DEVICE auto operator() (const XWindow &x_window_, const XResidualWindow &x_residual_window_, const XBiasWindow &x_bias_window_, const GammaWindow &gamma_window_, const BetaWindow &beta_window_, YWindow &y_window, const YResidualWindow &y_residual_window_, MeanWindow &mean_window, InvStdWindow &inv_std_window, const SmoothScaleWindow &, YScaleWindow &, ComputeDataType epsilon, ck_tile::index_t row_size, void *smem, Epilogue) const

Static Public Member Functions

static CK_TILE_HOST_DEVICE constexpr index_t GetSmemSize ()

Static Public Attributes

static constexpr bool kHasGamma = !std::is_same_v<GammaDataType, ck_tile::null_type>
static constexpr bool kHasBeta = !std::is_same_v<BetaDataType, ck_tile::null_type>
static constexpr bool kSaveMean = Problem::Traits::kSaveMeanInvStd
static constexpr bool kSaveInvStd = Problem::Traits::kSaveMeanInvStd
static constexpr bool kNeedCrossWarpSync = Problem::kNeedCrossWarpSync
static constexpr bool kPadM = false
static constexpr bool kPadN = Problem::Traits::kPadN
static constexpr bool kFastFDiv = Problem::Traits::kFastFDiv
static constexpr bool kWelford = Problem::Traits::kWelford
static constexpr auto kXbias = Problem::Traits::kXbias
static constexpr auto kFusedAdd = Problem::Traits::kFusedAdd
static constexpr auto kFusedQuant = Problem::Traits::kFusedQuant
static constexpr const char * name

Member Typedef Documentation

◆ BetaDataType

template<typename Problem_, typename Policy_ = Layernorm2dFwdPipelineDefaultPolicy>
using ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ >::BetaDataType = ck_tile::remove_cvref_t<typename Problem::BetaDataType>

◆ ComputeDataType

template<typename Problem_, typename Policy_ = Layernorm2dFwdPipelineDefaultPolicy>
using ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ >::ComputeDataType = ck_tile::remove_cvref_t<typename Problem::ComputeDataType>

◆ GammaDataType

template<typename Problem_, typename Policy_ = Layernorm2dFwdPipelineDefaultPolicy>
using ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ >::GammaDataType = ck_tile::remove_cvref_t<typename Problem::GammaDataType>

◆ InvStdDataType

template<typename Problem_, typename Policy_ = Layernorm2dFwdPipelineDefaultPolicy>
using ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ >::InvStdDataType = ck_tile::remove_cvref_t<typename Problem::InvStdDataType>

◆ MeanDataType

template<typename Problem_, typename Policy_ = Layernorm2dFwdPipelineDefaultPolicy>
using ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ >::MeanDataType = ck_tile::remove_cvref_t<typename Problem::MeanDataType>

◆ Policy

template<typename Problem_, typename Policy_ = Layernorm2dFwdPipelineDefaultPolicy>
using ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ >::Policy = ck_tile::remove_cvref_t<Policy_>

◆ Problem

template<typename Problem_, typename Policy_ = Layernorm2dFwdPipelineDefaultPolicy>
using ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ >::Problem = ck_tile::remove_cvref_t<Problem_>

◆ XBiasDataType

template<typename Problem_, typename Policy_ = Layernorm2dFwdPipelineDefaultPolicy>
using ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ >::XBiasDataType = ck_tile::remove_cvref_t<typename Problem::XBiasDataType>

◆ XDataType

template<typename Problem_, typename Policy_ = Layernorm2dFwdPipelineDefaultPolicy>
using ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ >::XDataType = ck_tile::remove_cvref_t<typename Problem::XDataType>

◆ XResidualDataType

template<typename Problem_, typename Policy_ = Layernorm2dFwdPipelineDefaultPolicy>
using ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ >::XResidualDataType = XDataType

◆ YDataType

template<typename Problem_, typename Policy_ = Layernorm2dFwdPipelineDefaultPolicy>
using ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ >::YDataType = ck_tile::remove_cvref_t<typename Problem::YDataType>

◆ YResidualDataType

template<typename Problem_, typename Policy_ = Layernorm2dFwdPipelineDefaultPolicy>
using ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ >::YResidualDataType = XDataType

Member Function Documentation

◆ GetSmemSize()

template<typename Problem_, typename Policy_ = Layernorm2dFwdPipelineDefaultPolicy>
CK_TILE_HOST_DEVICE constexpr index_t ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ >::GetSmemSize ( )
inlinestaticconstexpr

◆ operator()()

template<typename Problem_, typename Policy_ = Layernorm2dFwdPipelineDefaultPolicy>
template<typename XWindow, typename XResidualWindow, typename XBiasWindow, typename GammaWindow, typename BetaWindow, typename YWindow, typename YResidualWindow, typename MeanWindow, typename InvStdWindow, typename SmoothScaleWindow, typename YScaleWindow, typename Epilogue>
CK_TILE_DEVICE auto ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ >::operator() ( const XWindow & x_window_,
const XResidualWindow & x_residual_window_,
const XBiasWindow & x_bias_window_,
const GammaWindow & gamma_window_,
const BetaWindow & beta_window_,
YWindow & y_window,
const YResidualWindow & y_residual_window_,
MeanWindow & mean_window,
InvStdWindow & inv_std_window,
const SmoothScaleWindow & ,
YScaleWindow & ,
ComputeDataType epsilon,
ck_tile::index_t row_size,
void * smem,
Epilogue  ) const
inline

Member Data Documentation

◆ kFastFDiv

template<typename Problem_, typename Policy_ = Layernorm2dFwdPipelineDefaultPolicy>
bool ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ >::kFastFDiv = Problem::Traits::kFastFDiv
staticconstexpr

◆ kFusedAdd

template<typename Problem_, typename Policy_ = Layernorm2dFwdPipelineDefaultPolicy>
auto ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ >::kFusedAdd = Problem::Traits::kFusedAdd
staticconstexpr

◆ kFusedQuant

template<typename Problem_, typename Policy_ = Layernorm2dFwdPipelineDefaultPolicy>
auto ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ >::kFusedQuant = Problem::Traits::kFusedQuant
staticconstexpr

◆ kHasBeta

template<typename Problem_, typename Policy_ = Layernorm2dFwdPipelineDefaultPolicy>
bool ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ >::kHasBeta = !std::is_same_v<BetaDataType, ck_tile::null_type>
staticconstexpr

◆ kHasGamma

template<typename Problem_, typename Policy_ = Layernorm2dFwdPipelineDefaultPolicy>
bool ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ >::kHasGamma = !std::is_same_v<GammaDataType, ck_tile::null_type>
staticconstexpr

◆ kNeedCrossWarpSync

template<typename Problem_, typename Policy_ = Layernorm2dFwdPipelineDefaultPolicy>
bool ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ >::kNeedCrossWarpSync = Problem::kNeedCrossWarpSync
staticconstexpr

◆ kPadM

template<typename Problem_, typename Policy_ = Layernorm2dFwdPipelineDefaultPolicy>
bool ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ >::kPadM = false
staticconstexpr

◆ kPadN

template<typename Problem_, typename Policy_ = Layernorm2dFwdPipelineDefaultPolicy>
bool ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ >::kPadN = Problem::Traits::kPadN
staticconstexpr

◆ kSaveInvStd

template<typename Problem_, typename Policy_ = Layernorm2dFwdPipelineDefaultPolicy>
bool ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ >::kSaveInvStd = Problem::Traits::kSaveMeanInvStd
staticconstexpr

◆ kSaveMean

template<typename Problem_, typename Policy_ = Layernorm2dFwdPipelineDefaultPolicy>
bool ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ >::kSaveMean = Problem::Traits::kSaveMeanInvStd
staticconstexpr

◆ kWelford

template<typename Problem_, typename Policy_ = Layernorm2dFwdPipelineDefaultPolicy>
bool ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ >::kWelford = Problem::Traits::kWelford
staticconstexpr

◆ kXbias

template<typename Problem_, typename Policy_ = Layernorm2dFwdPipelineDefaultPolicy>
auto ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ >::kXbias = Problem::Traits::kXbias
staticconstexpr

◆ name

template<typename Problem_, typename Policy_ = Layernorm2dFwdPipelineDefaultPolicy>
const char* ck_tile::Layernorm2dFwdPipelineTwoPass< Problem_, Policy_ >::name
staticconstexpr
Initial value:
= []() {
if constexpr(kNeedCrossWarpSync)
return "bpr_2p";
else
return "wpr_2p";
}()
ck_tile::AddRmsnorm2dRdquantFwdPipelineOnePass::kNeedCrossWarpSync
static constexpr bool kNeedCrossWarpSync
Definition add_rmsnorm2d_rdquant_fwd_pipeline_one_pass.hpp:30
The documentation for this struct was generated from the following file: