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< ck_tile::bf16_t >
 
struct  t2s< ck_tile::bf8_t >
 
struct  t2s< ck_tile::fp16_t >
 
struct  t2s< ck_tile::fp8_t >
 
struct  t2s< ck_tile::int8_t >
 
struct  t2s< float >
 

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 XResidualDataType = XDataType
 
using YResidualDataType = XDataType
 
using Hargs = Rmsnorm2dFwdHostArgs
 

Public Member Functions

CK_TILE_DEVICE void operator() (Kargs kargs) const
 

Static Public Member Functions

static constexpr CK_TILE_HOST Kargs MakeKargs (const Hargs &hargs)
 
static constexpr CK_TILE_HOST auto GridSize (const Hargs &hargs)
 
static constexpr CK_TILE_HOST auto BlockSize ()
 
static constexpr CK_TILE_HOST_DEVICE 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 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 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 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>

◆ 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_ >
static constexpr CK_TILE_HOST auto ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::BlockSize ( )
inlinestaticconstexpr

◆ GetName()

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

◆ GetSmemSize()

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

◆ GridSize()

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

◆ MakeKargs()

template<typename Pipeline_ , typename Epilogue_ >
static constexpr CK_TILE_HOST 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_ >
constexpr index_t ck_tile::Rmsnorm2dFwd< Pipeline_, Epilogue_ >::Block_M = Problem::BlockShape::Block_M
staticconstexpr

◆ Block_N

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

◆ I0

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

◆ I1

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

◆ kFusedAdd

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

◆ kFusedQuant

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

◆ kHasGamma

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

◆ kPadM

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

◆ kPadN

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

◆ kSaveInvRms

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

◆ kTwoPass

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

◆ Repeat_N

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

◆ ThreadPerWarp_N

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

◆ Vector_N

template<typename Pipeline_ , typename Epilogue_ >
constexpr 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:
  • /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-6.4.3/include/ck_tile/ops/rmsnorm2d/kernel/rmsnorm2d_fwd_kernel.hpp