SmoothquantPipelineOnePass< Problem_, Policy_ > Struct Template Reference

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

Composable Kernel: ck_tile::SmoothquantPipelineOnePass< Problem_, Policy_ > Struct Template Reference
ck_tile::SmoothquantPipelineOnePass< Problem_, Policy_ > Struct Template Reference

#include <smoothquant_pipeline_one_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, void *smem) const
 

Static Public Member Functions

static constexpr CK_TILE_HOST_DEVICE 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::SmoothquantPipelineOnePass< Problem_, Policy_ >::ComputeDataType = ck_tile::remove_cvref_t<typename Problem::ComputeDataType>

◆ Policy

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

◆ Problem

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

◆ QYDataType

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

◆ SmoothScaleDataType

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

◆ XDataType

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

◆ YScaleDataType

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

Member Function Documentation

◆ GetSmemSize()

template<typename Problem_ , typename Policy_ = SmoothquantPipelineDefaultPolicy>
static constexpr CK_TILE_HOST_DEVICE index_t ck_tile::SmoothquantPipelineOnePass< 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::SmoothquantPipelineOnePass< Problem_, Policy_ >::operator() ( const XWindow &  x_window_,
const SmoothScaleWindow &  smscale_window_,
YScaleWindow &  yscale_window,
QYWindow &  qy_window,
ck_tile::index_t  ,
void *  smem 
) const
inline

Member Data Documentation

◆ kNeedCrossWarpSync

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

◆ kPadM

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

◆ kPadN

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

◆ name

template<typename Problem_ , typename Policy_ = SmoothquantPipelineDefaultPolicy>
constexpr const char* ck_tile::SmoothquantPipelineOnePass< Problem_, Policy_ >::name
staticconstexpr
Initial value:
= []() {
if constexpr(kNeedCrossWarpSync)
return "bpr_op";
else
return "wpr_op";
}()
static constexpr bool kNeedCrossWarpSync
Definition: smoothquant_pipeline_one_pass.hpp:25

◆ UseMax3

template<typename Problem_ , typename Policy_ = SmoothquantPipelineDefaultPolicy>
constexpr bool ck_tile::SmoothquantPipelineOnePass< Problem_, Policy_ >::UseMax3 = true
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/smoothquant/pipeline/smoothquant_pipeline_one_pass.hpp