Rmsnorm2dFwd< Pipeline_, Epilogue_ > Struct Template Reference

Rmsnorm2dFwd&lt; Pipeline_, Epilogue_ &gt; Struct Template Reference#

Composable Kernel: ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ > Struct Template Reference
Classes | Public Types | Public Member Functions | Static Public Member Functions | Static Public Attributes | List of all members
ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ > Struct Template Reference

#include <rmsnorm2d_fwd_kernel.hpp>

Classes

struct  Kargs
struct  t2s
struct  t2s< float >
struct  t2s< ck_tile::fp16_t >
struct  t2s< ck_tile::bf16_t >
struct  t2s< ck_tile::fp8_t >
struct  t2s< ck_tile::bf8_t >
struct  t2s< ck_tile::int8_t >

Public Types

using Pipeline = remove_cvref_t<Pipeline_>
using Epilogue = remove_cvref_t<Epilogue_>
using Problem = typename Pipeline::Problem
using XDataType = remove_cvref_t<typename Problem::XDataType>
using GammaDataType = remove_cvref_t<typename Problem::GammaDataType>
using ComputeDataType = remove_cvref_t<typename Problem::ComputeDataType>
using YDataType = remove_cvref_t<typename Problem::YDataType>
using InvRmsDataType = remove_cvref_t<typename Problem::InvRmsDataType>
using SmoothScaleDataType = remove_cvref_t<typename Problem::SmoothScaleDataType>
using YScaleDataType = remove_cvref_t<typename Problem::YScaleDataType>
using UnquantYDataType = remove_cvref_t<typename Problem::UnquantYDataType>
using XResidualDataType = XDataType
using YResidualDataType = XDataType
using Hargs = Rmsnorm2dFwdHostArgs

Public Member Functions

CK_TILE_DEVICE void operator() (Kargs kargs) const

Static Public Member Functions

static CK_TILE_HOST constexpr Kargs MakeKargs (const Hargs &hargs)
static CK_TILE_HOST constexpr auto GridSize (const Hargs &hargs)
static CK_TILE_HOST constexpr auto BlockSize ()
static CK_TILE_HOST_DEVICE constexpr index_t GetSmemSize ()
static CK_TILE_HOST std::string GetName ()

Static Public Attributes

static constexpr bool kHasGamma = !std::is_same_v<GammaDataType, null_type>
static constexpr bool kSaveInvRms = Problem::Traits::kSaveInvRms
static constexpr bool kSaveUnquant = Problem::Traits::kSaveUnquant
static constexpr index_t Block_M = Problem::BlockShape::Block_M
static constexpr index_t Block_N = Problem::BlockShape::Block_N
static constexpr bool kPadM = false
static constexpr bool kPadN = Problem::Traits::kPadN
static constexpr bool kTwoPass = Problem::Traits::kTwoPass
static constexpr auto kFusedAdd = Problem::Traits::kFusedAdd
static constexpr auto kFusedQuant = Problem::Traits::kFusedQuant
static constexpr auto kUseModelSensitiveRMSNorm = Problem::Traits::kUseModelSensitiveRMSNorm
static constexpr index_t ThreadPerWarp_N = Problem::BlockShape::ThreadPerWarp_N
static constexpr index_t Vector_N = Problem::BlockShape::Vector_N
static constexpr index_t Repeat_N = Problem::BlockShape::Repeat_N
static constexpr index_t kBlockSize = Problem::BlockShape::BlockSize
static constexpr auto I0 = number<0>{}
static constexpr auto I1 = number<1>{}

Member Typedef Documentation

◆ ComputeDataType

template<typename Pipeline_, typename Epilogue_>
using ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::ComputeDataType = remove_cvref_t<typename Problem::ComputeDataType>

◆ Epilogue

template<typename Pipeline_, typename Epilogue_>
using ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::Epilogue = remove_cvref_t<Epilogue_>

◆ GammaDataType

template<typename Pipeline_, typename Epilogue_>
using ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::GammaDataType = remove_cvref_t<typename Problem::GammaDataType>

◆ Hargs

template<typename Pipeline_, typename Epilogue_>
using ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::Hargs = Rmsnorm2dFwdHostArgs

◆ InvRmsDataType

template<typename Pipeline_, typename Epilogue_>
using ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::InvRmsDataType = remove_cvref_t<typename Problem::InvRmsDataType>

◆ Pipeline

template<typename Pipeline_, typename Epilogue_>
using ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::Pipeline = remove_cvref_t<Pipeline_>

◆ Problem

template<typename Pipeline_, typename Epilogue_>
using ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::Problem = typename Pipeline::Problem

◆ SmoothScaleDataType

template<typename Pipeline_, typename Epilogue_>
using ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::SmoothScaleDataType = remove_cvref_t<typename Problem::SmoothScaleDataType>

◆ UnquantYDataType

template<typename Pipeline_, typename Epilogue_>
using ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::UnquantYDataType = remove_cvref_t<typename Problem::UnquantYDataType>

◆ XDataType

template<typename Pipeline_, typename Epilogue_>
using ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::XDataType = remove_cvref_t<typename Problem::XDataType>

◆ XResidualDataType

template<typename Pipeline_, typename Epilogue_>
using ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::XResidualDataType = XDataType

◆ YDataType

template<typename Pipeline_, typename Epilogue_>
using ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::YDataType = remove_cvref_t<typename Problem::YDataType>

◆ YResidualDataType

template<typename Pipeline_, typename Epilogue_>
using ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::YResidualDataType = XDataType

◆ YScaleDataType

template<typename Pipeline_, typename Epilogue_>
using ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::YScaleDataType = remove_cvref_t<typename Problem::YScaleDataType>

Member Function Documentation

◆ BlockSize()

template<typename Pipeline_, typename Epilogue_>
CK_TILE_HOST constexpr auto ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::BlockSize ( )
inlinestaticconstexpr

◆ GetName()

template<typename Pipeline_, typename Epilogue_>
CK_TILE_HOST std::string ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::GetName ( )
inlinestatic

◆ GetSmemSize()

template<typename Pipeline_, typename Epilogue_>
CK_TILE_HOST_DEVICE constexpr index_t ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::GetSmemSize ( )
inlinestaticconstexpr

◆ GridSize()

template<typename Pipeline_, typename Epilogue_>
CK_TILE_HOST constexpr auto ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::GridSize ( const Hargs & hargs)
inlinestaticconstexpr

◆ MakeKargs()

template<typename Pipeline_, typename Epilogue_>
CK_TILE_HOST constexpr Kargs ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::MakeKargs ( const Hargs & hargs)
inlinestaticconstexpr

◆ operator()()

template<typename Pipeline_, typename Epilogue_>
CK_TILE_DEVICE void ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::operator() ( Kargs kargs) const
inline

Member Data Documentation

◆ Block_M

template<typename Pipeline_, typename Epilogue_>
index_t ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::Block_M = Problem::BlockShape::Block_M
staticconstexpr

◆ Block_N

template<typename Pipeline_, typename Epilogue_>
index_t ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::Block_N = Problem::BlockShape::Block_N
staticconstexpr

◆ I0

template<typename Pipeline_, typename Epilogue_>
auto ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::I0 = number<0>{}
staticconstexpr

◆ I1

template<typename Pipeline_, typename Epilogue_>
auto ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::I1 = number<1>{}
staticconstexpr

◆ kBlockSize

template<typename Pipeline_, typename Epilogue_>
index_t ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::kBlockSize = Problem::BlockShape::BlockSize
staticconstexpr

◆ kFusedAdd

template<typename Pipeline_, typename Epilogue_>
auto ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::kFusedAdd = Problem::Traits::kFusedAdd
staticconstexpr

◆ kFusedQuant

template<typename Pipeline_, typename Epilogue_>
auto ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::kFusedQuant = Problem::Traits::kFusedQuant
staticconstexpr

◆ kHasGamma

template<typename Pipeline_, typename Epilogue_>
bool ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::kHasGamma = !std::is_same_v<GammaDataType, null_type>
staticconstexpr

◆ kPadM

template<typename Pipeline_, typename Epilogue_>
bool ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::kPadM = false
staticconstexpr

◆ kPadN

template<typename Pipeline_, typename Epilogue_>
bool ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::kPadN = Problem::Traits::kPadN
staticconstexpr

◆ kSaveInvRms

template<typename Pipeline_, typename Epilogue_>
bool ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::kSaveInvRms = Problem::Traits::kSaveInvRms
staticconstexpr

◆ kSaveUnquant

template<typename Pipeline_, typename Epilogue_>
bool ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::kSaveUnquant = Problem::Traits::kSaveUnquant
staticconstexpr

◆ kTwoPass

template<typename Pipeline_, typename Epilogue_>
bool ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::kTwoPass = Problem::Traits::kTwoPass
staticconstexpr

◆ kUseModelSensitiveRMSNorm

template<typename Pipeline_, typename Epilogue_>
auto ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::kUseModelSensitiveRMSNorm = Problem::Traits::kUseModelSensitiveRMSNorm
staticconstexpr

◆ Repeat_N

template<typename Pipeline_, typename Epilogue_>
index_t ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::Repeat_N = Problem::BlockShape::Repeat_N
staticconstexpr

◆ ThreadPerWarp_N

template<typename Pipeline_, typename Epilogue_>
index_t ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::ThreadPerWarp_N = Problem::BlockShape::ThreadPerWarp_N
staticconstexpr

◆ Vector_N

template<typename Pipeline_, typename Epilogue_>
index_t ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::Vector_N = Problem::BlockShape::Vector_N
staticconstexpr
The documentation for this struct was generated from the following file: