/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/layernorm2d/kernel/layernorm2d_fwd_kernel.hpp Source File#
layernorm2d_fwd_kernel.hpp
Go to the documentation of this file.
160 if (kXbias != Layernorm2dXBiasEnum::NO_BIAS) n += _SS_("_") + Layernorm2dXBiasEnumName<kXbias>::name;
161 if (kFusedAdd != Layernorm2dFusedAddEnum::NO_ADD) n += _SS_("_") + Layernorm2dFusedAddEnumName<kFusedAdd>::name;
162 if (kFusedQuant != Layernorm2dFusedQuantEnum::NO_SWEEP) n += _SS_("_") + Layernorm2dFusedQuantEnumName<kFusedQuant>::name;
184 _TS_(S_::Block_M) + "x" + _TS_(S_::Block_N) + "_" + _TS_(S_::WarpPerBlock_M) + "x" + _TS_(S_::WarpPerBlock_N) + "_" +
185 _TS_(S_::Warp_M) + "x" + _TS_(S_::Warp_N) + "_" + _TS_(S_::Vector_M) + "x" + _TS_(S_::Vector_N) + "_" +
#define _TS_
#define _SS_
Definition: cluster_descriptor.hpp:13
constexpr CK_TILE_DEVICE auto make_null_tile_window(const WindowLengths &window_lengths)
Definition: null_tile_window.hpp:66
constexpr CK_TILE_HOST_DEVICE auto integer_divide_ceil(X x, Y y)
Definition: math.hpp:149
@ SMOOTH_DYNAMIC_QUANT
@ DYNAMIC_QUANT
@ ADD_BIAS
constexpr CK_TILE_HOST_DEVICE auto pad_tensor_view(const TensorView &tensor_view, const TileLengths &tile_lengths, DoPads)
Definition: tensor_view.hpp:530
remove_cv_t< std::remove_reference_t< T > > remove_cvref_t
Definition: type_traits.hpp:21
@ PRE_ADD_STORE
constexpr CK_TILE_DEVICE auto make_tile_window(null_tensor_view, const WindowLengths &window_lengths, const multi_index< WindowLengths::size()> &, Ts &&...)
Definition: null_tile_window.hpp:75
constexpr CK_TILE_HOST_DEVICE auto make_tuple(Xs &&... xs)
Definition: tuple.hpp:360
Definition: layernorm2d_fwd_traits.hpp:33
Definition: layernorm2d_fwd_traits.hpp:47
Definition: layernorm2d_fwd_kernel.hpp:84
index_t yr_stride
Definition: layernorm2d_fwd_kernel.hpp:106
const void * p_x_bias
Definition: layernorm2d_fwd_kernel.hpp:88
void * p_y_residual
Definition: layernorm2d_fwd_kernel.hpp:93
const void * p_gamma
Definition: layernorm2d_fwd_kernel.hpp:89
const void * p_sm_scale
Definition: layernorm2d_fwd_kernel.hpp:87
const void * p_x_residual
Definition: layernorm2d_fwd_kernel.hpp:86
index_t xr_stride
Definition: layernorm2d_fwd_kernel.hpp:104
Definition: layernorm2d_fwd_kernel.hpp:140
Definition: layernorm2d_fwd_kernel.hpp:14
index_t xr_stride
Definition: layernorm2d_fwd_kernel.hpp:33
void * p_y_residual
Definition: layernorm2d_fwd_kernel.hpp:23
const void * p_x_bias
Definition: layernorm2d_fwd_kernel.hpp:18
const void * p_gamma
Definition: layernorm2d_fwd_kernel.hpp:19
const void * p_x_residual
Definition: layernorm2d_fwd_kernel.hpp:16
const void * p_sm_scale
Definition: layernorm2d_fwd_kernel.hpp:17
index_t yr_stride
Definition: layernorm2d_fwd_kernel.hpp:35
Definition: layernorm2d_fwd_kernel.hpp:41
typename Pipeline::Problem Problem
Definition: layernorm2d_fwd_kernel.hpp:44
static constexpr CK_TILE_HOST auto GridSize(const Hargs &hargs)
Definition: layernorm2d_fwd_kernel.hpp:132
remove_cvref_t< typename Problem::BetaDataType > BetaDataType
Definition: layernorm2d_fwd_kernel.hpp:49
remove_cvref_t< Pipeline_ > Pipeline
Definition: layernorm2d_fwd_kernel.hpp:42
remove_cvref_t< typename Problem::XDataType > XDataType
Definition: layernorm2d_fwd_kernel.hpp:46
static constexpr bool kHasBeta
Definition: layernorm2d_fwd_kernel.hpp:62
static CK_TILE_HOST std::string GetName()
Definition: layernorm2d_fwd_kernel.hpp:152
static constexpr auto kFusedAdd
Definition: layernorm2d_fwd_kernel.hpp:73
static constexpr index_t Repeat_N
Definition: layernorm2d_fwd_kernel.hpp:78
static constexpr CK_TILE_HOST auto BlockSize()
Definition: layernorm2d_fwd_kernel.hpp:137
static constexpr CK_TILE_HOST_DEVICE index_t GetSmemSize()
Definition: layernorm2d_fwd_kernel.hpp:150
static constexpr CK_TILE_HOST Kargs MakeKargs(const Hargs &hargs)
Definition: layernorm2d_fwd_kernel.hpp:110
static constexpr index_t Block_N
Definition: layernorm2d_fwd_kernel.hpp:68
remove_cvref_t< typename Problem::XBiasDataType > XBiasDataType
Definition: layernorm2d_fwd_kernel.hpp:47
static constexpr bool kHasGamma
Definition: layernorm2d_fwd_kernel.hpp:61
static constexpr index_t ThreadPerWarp_N
Definition: layernorm2d_fwd_kernel.hpp:76
static constexpr bool kSaveMeanInvStd
Definition: layernorm2d_fwd_kernel.hpp:63
CK_TILE_DEVICE void operator()(Kargs kargs) const
Definition: layernorm2d_fwd_kernel.hpp:192
remove_cvref_t< typename Problem::ComputeDataType > ComputeDataType
Definition: layernorm2d_fwd_kernel.hpp:50
remove_cvref_t< typename Problem::YScaleDataType > YScaleDataType
Definition: layernorm2d_fwd_kernel.hpp:55
static constexpr index_t Vector_N
Definition: layernorm2d_fwd_kernel.hpp:77
static constexpr bool kTwoPass
Definition: layernorm2d_fwd_kernel.hpp:71
XDataType YResidualDataType
Definition: layernorm2d_fwd_kernel.hpp:59
static constexpr auto kFusedQuant
Definition: layernorm2d_fwd_kernel.hpp:74
static constexpr bool kSaveMean
Definition: layernorm2d_fwd_kernel.hpp:64
remove_cvref_t< typename Problem::YDataType > YDataType
Definition: layernorm2d_fwd_kernel.hpp:51
static constexpr auto kXbias
Definition: layernorm2d_fwd_kernel.hpp:72
static constexpr index_t Block_M
Definition: layernorm2d_fwd_kernel.hpp:67
XDataType XResidualDataType
Definition: layernorm2d_fwd_kernel.hpp:58
remove_cvref_t< typename Problem::InvStdDataType > InvStdDataType
Definition: layernorm2d_fwd_kernel.hpp:53
remove_cvref_t< typename Problem::MeanDataType > MeanDataType
Definition: layernorm2d_fwd_kernel.hpp:52
static constexpr bool kSaveInvStd
Definition: layernorm2d_fwd_kernel.hpp:65
static constexpr index_t kBlockSize
Definition: layernorm2d_fwd_kernel.hpp:79
remove_cvref_t< typename Problem::SmoothScaleDataType > SmoothScaleDataType
Definition: layernorm2d_fwd_kernel.hpp:54
remove_cvref_t< Epilogue_ > Epilogue
Definition: layernorm2d_fwd_kernel.hpp:43
remove_cvref_t< typename Problem::GammaDataType > GammaDataType
Definition: layernorm2d_fwd_kernel.hpp:48
Definition: layernorm2d_fwd_traits.hpp:18
Definition: integral_constant.hpp:13
Definition: sequence.hpp:49