AddRmsnorm2dRdquantFwdPipelineOnePass< Problem_, Policy_ > Struct Template Reference

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

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

#include <add_rmsnorm2d_rdquant_fwd_pipeline_one_pass.hpp>

Public Types

using Problem = ck_tile::remove_cvref_t<Problem_>
using Policy = ck_tile::remove_cvref_t<Policy_>
using ADataType = ck_tile::remove_cvref_t<typename Problem::ADataType>
using BDataType = ck_tile::remove_cvref_t<typename Problem::BDataType>
using GammaDataType = ck_tile::remove_cvref_t<typename Problem::GammaDataType>
using ComputeDataType = ck_tile::remove_cvref_t<typename Problem::ComputeDataType>
using XDataType = ck_tile::remove_cvref_t<typename Problem::XDataType>
using YScaleDataType = ck_tile::remove_cvref_t<typename Problem::YScaleDataType>
using QYDataType = ck_tile::remove_cvref_t<typename Problem::QYDataType>

Public Member Functions

template<typename AWindow, typename BWindow, typename GammaWindow, typename XWindow, typename YScaleWindow, typename QYWindow>
CK_TILE_DEVICE auto operator() (const AWindow &a_window_, const BWindow &b_window_, const GammaWindow &gamma_window_, XWindow &x_window, YScaleWindow &yscale_window, QYWindow &qy_window, ComputeDataType epsilon, ck_tile::index_t row_size, void *smem) 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 kSaveX = Problem::kSaveX
static constexpr bool kNeedCrossWarpSync = Problem::kNeedCrossWarpSync
static constexpr bool kPadM = false
static constexpr bool kPadN = Problem::kPadN
static constexpr bool UseMax3 = true
static constexpr const char * name

Member Typedef Documentation

◆ ADataType

template<typename Problem_, typename Policy_ = AddRmsnorm2dRdquantFwdPipelineDefaultPolicy>
using ck_tile::AddRmsnorm2dRdquantFwdPipelineOnePass< Problem_, Policy_ >::ADataType = ck_tile::remove_cvref_t<typename Problem::ADataType>

◆ BDataType

template<typename Problem_, typename Policy_ = AddRmsnorm2dRdquantFwdPipelineDefaultPolicy>
using ck_tile::AddRmsnorm2dRdquantFwdPipelineOnePass< Problem_, Policy_ >::BDataType = ck_tile::remove_cvref_t<typename Problem::BDataType>

◆ ComputeDataType

template<typename Problem_, typename Policy_ = AddRmsnorm2dRdquantFwdPipelineDefaultPolicy>
using ck_tile::AddRmsnorm2dRdquantFwdPipelineOnePass< Problem_, Policy_ >::ComputeDataType = ck_tile::remove_cvref_t<typename Problem::ComputeDataType>

◆ GammaDataType

template<typename Problem_, typename Policy_ = AddRmsnorm2dRdquantFwdPipelineDefaultPolicy>
using ck_tile::AddRmsnorm2dRdquantFwdPipelineOnePass< Problem_, Policy_ >::GammaDataType = ck_tile::remove_cvref_t<typename Problem::GammaDataType>

◆ Policy

template<typename Problem_, typename Policy_ = AddRmsnorm2dRdquantFwdPipelineDefaultPolicy>
using ck_tile::AddRmsnorm2dRdquantFwdPipelineOnePass< Problem_, Policy_ >::Policy = ck_tile::remove_cvref_t<Policy_>

◆ Problem

template<typename Problem_, typename Policy_ = AddRmsnorm2dRdquantFwdPipelineDefaultPolicy>
using ck_tile::AddRmsnorm2dRdquantFwdPipelineOnePass< Problem_, Policy_ >::Problem = ck_tile::remove_cvref_t<Problem_>

◆ QYDataType

template<typename Problem_, typename Policy_ = AddRmsnorm2dRdquantFwdPipelineDefaultPolicy>
using ck_tile::AddRmsnorm2dRdquantFwdPipelineOnePass< Problem_, Policy_ >::QYDataType = ck_tile::remove_cvref_t<typename Problem::QYDataType>

◆ XDataType

template<typename Problem_, typename Policy_ = AddRmsnorm2dRdquantFwdPipelineDefaultPolicy>
using ck_tile::AddRmsnorm2dRdquantFwdPipelineOnePass< Problem_, Policy_ >::XDataType = ck_tile::remove_cvref_t<typename Problem::XDataType>

◆ YScaleDataType

template<typename Problem_, typename Policy_ = AddRmsnorm2dRdquantFwdPipelineDefaultPolicy>
using ck_tile::AddRmsnorm2dRdquantFwdPipelineOnePass< Problem_, Policy_ >::YScaleDataType = ck_tile::remove_cvref_t<typename Problem::YScaleDataType>

Member Function Documentation

◆ GetSmemSize()

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

◆ operator()()

template<typename Problem_, typename Policy_ = AddRmsnorm2dRdquantFwdPipelineDefaultPolicy>
template<typename AWindow, typename BWindow, typename GammaWindow, typename XWindow, typename YScaleWindow, typename QYWindow>
CK_TILE_DEVICE auto ck_tile::AddRmsnorm2dRdquantFwdPipelineOnePass< Problem_, Policy_ >::operator() ( const AWindow & a_window_,
const BWindow & b_window_,
const GammaWindow & gamma_window_,
XWindow & x_window,
YScaleWindow & yscale_window,
QYWindow & qy_window,
ComputeDataType epsilon,
ck_tile::index_t row_size,
void * smem ) const
inline

Member Data Documentation

◆ kHasGamma

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

◆ kNeedCrossWarpSync

template<typename Problem_, typename Policy_ = AddRmsnorm2dRdquantFwdPipelineDefaultPolicy>
bool ck_tile::AddRmsnorm2dRdquantFwdPipelineOnePass< Problem_, Policy_ >::kNeedCrossWarpSync = Problem::kNeedCrossWarpSync
staticconstexpr

◆ kPadM

template<typename Problem_, typename Policy_ = AddRmsnorm2dRdquantFwdPipelineDefaultPolicy>
bool ck_tile::AddRmsnorm2dRdquantFwdPipelineOnePass< Problem_, Policy_ >::kPadM = false
staticconstexpr

◆ kPadN

template<typename Problem_, typename Policy_ = AddRmsnorm2dRdquantFwdPipelineDefaultPolicy>
bool ck_tile::AddRmsnorm2dRdquantFwdPipelineOnePass< Problem_, Policy_ >::kPadN = Problem::kPadN
staticconstexpr

◆ kSaveX

template<typename Problem_, typename Policy_ = AddRmsnorm2dRdquantFwdPipelineDefaultPolicy>
bool ck_tile::AddRmsnorm2dRdquantFwdPipelineOnePass< Problem_, Policy_ >::kSaveX = Problem::kSaveX
staticconstexpr

◆ name

template<typename Problem_, typename Policy_ = AddRmsnorm2dRdquantFwdPipelineDefaultPolicy>
const char* ck_tile::AddRmsnorm2dRdquantFwdPipelineOnePass< Problem_, Policy_ >::name
staticconstexpr
Initial value:
= []() {
if constexpr(kNeedCrossWarpSync)
return "bpr_op";
else
return "wpr_op";
}()
ck_tile::AddRmsnorm2dRdquantFwdPipelineOnePass::kNeedCrossWarpSync
static constexpr bool kNeedCrossWarpSync
Definition add_rmsnorm2d_rdquant_fwd_pipeline_one_pass.hpp:30

◆ UseMax3

template<typename Problem_, typename Policy_ = AddRmsnorm2dRdquantFwdPipelineDefaultPolicy>
bool ck_tile::AddRmsnorm2dRdquantFwdPipelineOnePass< Problem_, Policy_ >::UseMax3 = true
staticconstexpr
The documentation for this struct was generated from the following file: