DynamicQuantEpilogue< Problem_, Policy_ > Struct Template Reference

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

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

#include <dynamic_quant_epilogue.hpp>

Public Types

using Problem = remove_cvref_t<Problem_>
using AccDataType = remove_cvref_t<typename Problem::AccDataType>
using SmoothScaleDataType = remove_cvref_t<typename Problem::SmoothScaleDataType>
using YScaleDataType = remove_cvref_t<typename Problem::YScaleDataType>
using ODataType = remove_cvref_t<typename Problem::ODataType>
using BlockShape = remove_cvref_t<typename Problem::BlockShape>

Public Member Functions

template<typename ODramWindowTmp, typename YScaleWindow, typename OAccTile>
CK_TILE_DEVICE auto Impl (ODramWindowTmp &o_dram_window_tmp, YScaleWindow &y_scale_window, const OAccTile &o_acc_tile, void *smem)
template<typename ODramWindowTmp, typename SmoothScaleWindow, typename YScaleWindow, typename OAccTile>
CK_TILE_DEVICE auto operator() (ODramWindowTmp &o_dram_window_tmp, const SmoothScaleWindow &sm_scale_window_, YScaleWindow &y_scale_window, const OAccTile &o_acc_tile, void *smem)
template<typename ODramWindowTmp, typename YScaleWindow, typename OAccTile>
CK_TILE_DEVICE auto operator() (ODramWindowTmp &o_dram_window_tmp, YScaleWindow &y_scale_window, const OAccTile &o_acc_tile, void *smem)

Static Public Member Functions

static CK_TILE_HOST_DEVICE constexpr auto GetBlockReduce2d ()
static CK_TILE_HOST_DEVICE constexpr auto GetBlockReduce2dSync ()
static CK_TILE_HOST_DEVICE constexpr auto GetBlockReduce2dCrossWarpSync ()
static CK_TILE_DEVICE constexpr auto MakeSmoothInputScaleTileDistribution ()
static CK_TILE_HOST_DEVICE constexpr index_t GetSmemSize ()

Static Public Attributes

static constexpr bool kPadM = Problem::Traits::kPadM
static constexpr bool kPadN = Problem::Traits::kPadN
static constexpr bool UseRawStore = Problem::Traits::UseRawStore
static constexpr bool UseMax3 = Problem::Traits::UseMax3

Member Typedef Documentation

◆ AccDataType

template<typename Problem_, typename Policy_ = void>
using ck_tile::DynamicQuantEpilogue< Problem_, Policy_ >::AccDataType = remove_cvref_t<typename Problem::AccDataType>

◆ BlockShape

template<typename Problem_, typename Policy_ = void>
using ck_tile::DynamicQuantEpilogue< Problem_, Policy_ >::BlockShape = remove_cvref_t<typename Problem::BlockShape>

◆ ODataType

template<typename Problem_, typename Policy_ = void>
using ck_tile::DynamicQuantEpilogue< Problem_, Policy_ >::ODataType = remove_cvref_t<typename Problem::ODataType>

◆ Problem

template<typename Problem_, typename Policy_ = void>
using ck_tile::DynamicQuantEpilogue< Problem_, Policy_ >::Problem = remove_cvref_t<Problem_>

◆ SmoothScaleDataType

template<typename Problem_, typename Policy_ = void>
using ck_tile::DynamicQuantEpilogue< Problem_, Policy_ >::SmoothScaleDataType = remove_cvref_t<typename Problem::SmoothScaleDataType>

◆ YScaleDataType

template<typename Problem_, typename Policy_ = void>
using ck_tile::DynamicQuantEpilogue< Problem_, Policy_ >::YScaleDataType = remove_cvref_t<typename Problem::YScaleDataType>

Member Function Documentation

◆ GetBlockReduce2d()

template<typename Problem_, typename Policy_ = void>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::DynamicQuantEpilogue< Problem_, Policy_ >::GetBlockReduce2d ( )
inlinestaticconstexpr

◆ GetBlockReduce2dCrossWarpSync()

template<typename Problem_, typename Policy_ = void>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::DynamicQuantEpilogue< Problem_, Policy_ >::GetBlockReduce2dCrossWarpSync ( )
inlinestaticconstexpr

◆ GetBlockReduce2dSync()

template<typename Problem_, typename Policy_ = void>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::DynamicQuantEpilogue< Problem_, Policy_ >::GetBlockReduce2dSync ( )
inlinestaticconstexpr

◆ GetSmemSize()

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

◆ Impl()

template<typename Problem_, typename Policy_ = void>
template<typename ODramWindowTmp, typename YScaleWindow, typename OAccTile>
CK_TILE_DEVICE auto ck_tile::DynamicQuantEpilogue< Problem_, Policy_ >::Impl ( ODramWindowTmp & o_dram_window_tmp,
YScaleWindow & y_scale_window,
const OAccTile & o_acc_tile,
void * smem )
inline

◆ MakeSmoothInputScaleTileDistribution()

template<typename Problem_, typename Policy_ = void>
CK_TILE_DEVICE constexpr auto ck_tile::DynamicQuantEpilogue< Problem_, Policy_ >::MakeSmoothInputScaleTileDistribution ( )
inlinestaticconstexpr

◆ operator()() [1/2]

template<typename Problem_, typename Policy_ = void>
template<typename ODramWindowTmp, typename SmoothScaleWindow, typename YScaleWindow, typename OAccTile>
CK_TILE_DEVICE auto ck_tile::DynamicQuantEpilogue< Problem_, Policy_ >::operator() ( ODramWindowTmp & o_dram_window_tmp,
const SmoothScaleWindow & sm_scale_window_,
YScaleWindow & y_scale_window,
const OAccTile & o_acc_tile,
void * smem )
inline

◆ operator()() [2/2]

template<typename Problem_, typename Policy_ = void>
template<typename ODramWindowTmp, typename YScaleWindow, typename OAccTile>
CK_TILE_DEVICE auto ck_tile::DynamicQuantEpilogue< Problem_, Policy_ >::operator() ( ODramWindowTmp & o_dram_window_tmp,
YScaleWindow & y_scale_window,
const OAccTile & o_acc_tile,
void * smem )
inline

Member Data Documentation

◆ kPadM

template<typename Problem_, typename Policy_ = void>
bool ck_tile::DynamicQuantEpilogue< Problem_, Policy_ >::kPadM = Problem::Traits::kPadM
staticconstexpr

◆ kPadN

template<typename Problem_, typename Policy_ = void>
bool ck_tile::DynamicQuantEpilogue< Problem_, Policy_ >::kPadN = Problem::Traits::kPadN
staticconstexpr

◆ UseMax3

template<typename Problem_, typename Policy_ = void>
bool ck_tile::DynamicQuantEpilogue< Problem_, Policy_ >::UseMax3 = Problem::Traits::UseMax3
staticconstexpr

◆ UseRawStore

template<typename Problem_, typename Policy_ = void>
bool ck_tile::DynamicQuantEpilogue< Problem_, Policy_ >::UseRawStore = Problem::Traits::UseRawStore
staticconstexpr
The documentation for this struct was generated from the following file: