SmoothquantPipelineTwoPass< Problem_, Policy_ > Struct Template Reference

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

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

#include <smoothquant_pipeline_two_pass.hpp>

Public Types

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

Public Member Functions

template<typename XWindow, typename SmoothScaleWindow, typename QYWindow, typename YScaleWindow>
CK_TILE_DEVICE auto operator() (const XWindow &x_window_, const SmoothScaleWindow &smscale_window_, YScaleWindow &yscale_window, QYWindow &qy_window, 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 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

◆ ComputeDataType

template<typename Problem_, typename Policy_ = SmoothquantPipelineDefaultPolicy>
using ck_tile::SmoothquantPipelineTwoPass< Problem_, Policy_ >::ComputeDataType = ck_tile::remove_cvref_t<typename Problem::ComputeDataType>

◆ Policy

template<typename Problem_, typename Policy_ = SmoothquantPipelineDefaultPolicy>
using ck_tile::SmoothquantPipelineTwoPass< Problem_, Policy_ >::Policy = ck_tile::remove_cvref_t<Policy_>

◆ Problem

template<typename Problem_, typename Policy_ = SmoothquantPipelineDefaultPolicy>
using ck_tile::SmoothquantPipelineTwoPass< Problem_, Policy_ >::Problem = ck_tile::remove_cvref_t<Problem_>

◆ QYDataType

template<typename Problem_, typename Policy_ = SmoothquantPipelineDefaultPolicy>
using ck_tile::SmoothquantPipelineTwoPass< Problem_, Policy_ >::QYDataType = ck_tile::remove_cvref_t<typename Problem::QYDataType>

◆ SmoothScaleDataType

template<typename Problem_, typename Policy_ = SmoothquantPipelineDefaultPolicy>
using ck_tile::SmoothquantPipelineTwoPass< Problem_, Policy_ >::SmoothScaleDataType = ck_tile::remove_cvref_t<typename Problem::SmoothScaleDataType>

◆ XDataType

template<typename Problem_, typename Policy_ = SmoothquantPipelineDefaultPolicy>
using ck_tile::SmoothquantPipelineTwoPass< Problem_, Policy_ >::XDataType = ck_tile::remove_cvref_t<typename Problem::XDataType>

◆ YScaleDataType

template<typename Problem_, typename Policy_ = SmoothquantPipelineDefaultPolicy>
using ck_tile::SmoothquantPipelineTwoPass< Problem_, Policy_ >::YScaleDataType = ck_tile::remove_cvref_t<typename Problem::YScaleDataType>

Member Function Documentation

◆ GetSmemSize()

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

◆ operator()()

template<typename Problem_, typename Policy_ = SmoothquantPipelineDefaultPolicy>
template<typename XWindow, typename SmoothScaleWindow, typename QYWindow, typename YScaleWindow>
CK_TILE_DEVICE auto ck_tile::SmoothquantPipelineTwoPass< Problem_, Policy_ >::operator() ( const XWindow & x_window_,
const SmoothScaleWindow & smscale_window_,
YScaleWindow & yscale_window,
QYWindow & qy_window,
ck_tile::index_t row_size,
void * smem ) const
inline

Member Data Documentation

◆ kNeedCrossWarpSync

template<typename Problem_, typename Policy_ = SmoothquantPipelineDefaultPolicy>
bool ck_tile::SmoothquantPipelineTwoPass< Problem_, Policy_ >::kNeedCrossWarpSync = Problem::kNeedCrossWarpSync
staticconstexpr

◆ kPadM

template<typename Problem_, typename Policy_ = SmoothquantPipelineDefaultPolicy>
bool ck_tile::SmoothquantPipelineTwoPass< Problem_, Policy_ >::kPadM = false
staticconstexpr

◆ kPadN

template<typename Problem_, typename Policy_ = SmoothquantPipelineDefaultPolicy>
bool ck_tile::SmoothquantPipelineTwoPass< Problem_, Policy_ >::kPadN = Problem::kPadN
staticconstexpr

◆ name

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

◆ UseMax3

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