GemmQuantPipelineProblemBase< ADataType_, AQDataType_, BDataType_, BQDataType_, CDataType_, BlockGemmShape_, Traits_, QuantGroupSize_, TransposeC_, ComputeDataType_, Scheduler_, HasHotLoop_, TailNum_ > Struct Template Reference

GemmQuantPipelineProblemBase&lt; ADataType_, AQDataType_, BDataType_, BQDataType_, CDataType_, BlockGemmShape_, Traits_, QuantGroupSize_, TransposeC_, ComputeDataType_, Scheduler_, HasHotLoop_, TailNum_ &gt; Struct Template Reference#

Composable Kernel: ck_tile::GemmQuantPipelineProblemBase< ADataType_, AQDataType_, BDataType_, BQDataType_, CDataType_, BlockGemmShape_, Traits_, QuantGroupSize_, TransposeC_, ComputeDataType_, Scheduler_, HasHotLoop_, TailNum_ > Struct Template Reference
ck_tile::GemmQuantPipelineProblemBase< ADataType_, AQDataType_, BDataType_, BQDataType_, CDataType_, BlockGemmShape_, Traits_, QuantGroupSize_, TransposeC_, ComputeDataType_, Scheduler_, HasHotLoop_, TailNum_ > Struct Template Reference

#include <gemm_quant_pipeline_problem.hpp>

Inheritance diagram for ck_tile::GemmQuantPipelineProblemBase< ADataType_, AQDataType_, BDataType_, BQDataType_, CDataType_, BlockGemmShape_, Traits_, QuantGroupSize_, TransposeC_, ComputeDataType_, Scheduler_, HasHotLoop_, TailNum_ >:
ck_tile::GemmPipelineProblemBase< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, BDataType_ >

Public Types

using Base
using Traits = typename Base::Traits
using AQDataType = remove_cvref_t<AQDataType_>
using BQDataType = remove_cvref_t<BQDataType_>
using BlockGemmShape = typename Base::BlockGemmShape
using QuantGroupSize = QuantGroupSize_
using AQLayout = remove_cvref_t<typename Traits::AQLayout>
using BQLayout = remove_cvref_t<typename Traits::BQLayout>
Public Types inherited from ck_tile::GemmPipelineProblemBase< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, BDataType_ >
using Traits
using AsDataType
using BsDataType
using CDataType
using BlockGemmShape
using AElementWise
using BElementWise
using AsLayout
using BsLayout
using CLayout
using ComputeDataTypeTuple
using AsLayoutTuple
using BsLayoutTuple
using AsDataTypeTuple
using BsDataTypeTuple
using ComputeDataType
using ADataType
using ALayout
using BDataType
using BLayout

Static Public Member Functions

static CK_TILE_HOST const std::string GetName ()
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentAQ ()
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentBQ ()
Static Public Member Functions inherited from ck_tile::GemmPipelineProblemBase< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, BDataType_ >
static CK_TILE_HOST const std::string GetName ()
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentA ()
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentB ()
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentC ()

Static Public Attributes

static constexpr bool TransposeC = TransposeC_
static constexpr bool PreshuffleB = Traits::PreshuffleB
static constexpr bool DoubleSmemBuffer = Traits::DoubleSmemBuffer
static constexpr auto Scheduler = Scheduler_
static constexpr auto HasHotLoop = HasHotLoop_
static constexpr auto TailNum = TailNum_
static constexpr index_t VectorSizeAQ
static constexpr index_t VectorSizeBQ = []() { return kPadK ? 1 : GetAlignmentBQ(); }()
static constexpr index_t kBlockSize
static constexpr bool kPadK
static constexpr bool kPadM
static constexpr bool kPadN
static constexpr index_t VectorLoadSize
Static Public Attributes inherited from ck_tile::GemmPipelineProblemBase< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, BDataType_ >
static constexpr bool FixedVectorSize
static constexpr bool ComputeDataTypeIsTuple
static constexpr bool ADataTypeIsTuple
static constexpr bool BDataTypeIsTuple
static constexpr bool ALayoutIsTuple
static constexpr bool BLayoutIsTuple
static constexpr bool TransposeC
static constexpr index_t NumWaveGroups
static constexpr bool UseStructuredSparsity
static constexpr index_t kBlockSize
static constexpr bool kPadM
static constexpr bool kPadN
static constexpr bool kPadK
static constexpr bool DoubleSmemBuffer
static constexpr auto Scheduler
static constexpr index_t VectorLoadSize
static constexpr bool Preshuffle
static constexpr index_t VectorSizeA
static constexpr index_t VectorSizeB
static constexpr index_t VectorSizeC

Member Typedef Documentation

◆ AQDataType

template<typename ADataType_, typename AQDataType_, typename BDataType_, typename BQDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, typename QuantGroupSize_, bool TransposeC_, typename ComputeDataType_ = BDataType_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full>
using ck_tile::GemmQuantPipelineProblemBase< ADataType_, AQDataType_, BDataType_, BQDataType_, CDataType_, BlockGemmShape_, Traits_, QuantGroupSize_, TransposeC_, ComputeDataType_, Scheduler_, HasHotLoop_, TailNum_ >::AQDataType = remove_cvref_t<AQDataType_>

◆ AQLayout

template<typename ADataType_, typename AQDataType_, typename BDataType_, typename BQDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, typename QuantGroupSize_, bool TransposeC_, typename ComputeDataType_ = BDataType_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full>
using ck_tile::GemmQuantPipelineProblemBase< ADataType_, AQDataType_, BDataType_, BQDataType_, CDataType_, BlockGemmShape_, Traits_, QuantGroupSize_, TransposeC_, ComputeDataType_, Scheduler_, HasHotLoop_, TailNum_ >::AQLayout = remove_cvref_t<typename Traits::AQLayout>

◆ Base

template<typename ADataType_, typename AQDataType_, typename BDataType_, typename BQDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, typename QuantGroupSize_, bool TransposeC_, typename ComputeDataType_ = BDataType_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full>
using ck_tile::GemmQuantPipelineProblemBase< ADataType_, AQDataType_, BDataType_, BQDataType_, CDataType_, BlockGemmShape_, Traits_, QuantGroupSize_, TransposeC_, ComputeDataType_, Scheduler_, HasHotLoop_, TailNum_ >::Base
Initial value:
BDataType_,
CDataType_,
BlockGemmShape_,
Traits_,
ComputeDataType_>
Definition gemm_pipeline_problem.hpp:25

◆ BlockGemmShape

template<typename ADataType_, typename AQDataType_, typename BDataType_, typename BQDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, typename QuantGroupSize_, bool TransposeC_, typename ComputeDataType_ = BDataType_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full>
using ck_tile::GemmQuantPipelineProblemBase< ADataType_, AQDataType_, BDataType_, BQDataType_, CDataType_, BlockGemmShape_, Traits_, QuantGroupSize_, TransposeC_, ComputeDataType_, Scheduler_, HasHotLoop_, TailNum_ >::BlockGemmShape = typename Base::BlockGemmShape

◆ BQDataType

template<typename ADataType_, typename AQDataType_, typename BDataType_, typename BQDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, typename QuantGroupSize_, bool TransposeC_, typename ComputeDataType_ = BDataType_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full>
using ck_tile::GemmQuantPipelineProblemBase< ADataType_, AQDataType_, BDataType_, BQDataType_, CDataType_, BlockGemmShape_, Traits_, QuantGroupSize_, TransposeC_, ComputeDataType_, Scheduler_, HasHotLoop_, TailNum_ >::BQDataType = remove_cvref_t<BQDataType_>

◆ BQLayout

template<typename ADataType_, typename AQDataType_, typename BDataType_, typename BQDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, typename QuantGroupSize_, bool TransposeC_, typename ComputeDataType_ = BDataType_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full>
using ck_tile::GemmQuantPipelineProblemBase< ADataType_, AQDataType_, BDataType_, BQDataType_, CDataType_, BlockGemmShape_, Traits_, QuantGroupSize_, TransposeC_, ComputeDataType_, Scheduler_, HasHotLoop_, TailNum_ >::BQLayout = remove_cvref_t<typename Traits::BQLayout>

◆ QuantGroupSize

template<typename ADataType_, typename AQDataType_, typename BDataType_, typename BQDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, typename QuantGroupSize_, bool TransposeC_, typename ComputeDataType_ = BDataType_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full>
using ck_tile::GemmQuantPipelineProblemBase< ADataType_, AQDataType_, BDataType_, BQDataType_, CDataType_, BlockGemmShape_, Traits_, QuantGroupSize_, TransposeC_, ComputeDataType_, Scheduler_, HasHotLoop_, TailNum_ >::QuantGroupSize = QuantGroupSize_

◆ Traits

template<typename ADataType_, typename AQDataType_, typename BDataType_, typename BQDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, typename QuantGroupSize_, bool TransposeC_, typename ComputeDataType_ = BDataType_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full>
using ck_tile::GemmQuantPipelineProblemBase< ADataType_, AQDataType_, BDataType_, BQDataType_, CDataType_, BlockGemmShape_, Traits_, QuantGroupSize_, TransposeC_, ComputeDataType_, Scheduler_, HasHotLoop_, TailNum_ >::Traits = typename Base::Traits

Member Function Documentation

◆ GetAlignmentAQ()

template<typename ADataType_, typename AQDataType_, typename BDataType_, typename BQDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, typename QuantGroupSize_, bool TransposeC_, typename ComputeDataType_ = BDataType_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::GemmQuantPipelineProblemBase< ADataType_, AQDataType_, BDataType_, BQDataType_, CDataType_, BlockGemmShape_, Traits_, QuantGroupSize_, TransposeC_, ComputeDataType_, Scheduler_, HasHotLoop_, TailNum_ >::GetAlignmentAQ ( )
inlinestaticconstexpr

◆ GetAlignmentBQ()

template<typename ADataType_, typename AQDataType_, typename BDataType_, typename BQDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, typename QuantGroupSize_, bool TransposeC_, typename ComputeDataType_ = BDataType_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::GemmQuantPipelineProblemBase< ADataType_, AQDataType_, BDataType_, BQDataType_, CDataType_, BlockGemmShape_, Traits_, QuantGroupSize_, TransposeC_, ComputeDataType_, Scheduler_, HasHotLoop_, TailNum_ >::GetAlignmentBQ ( )
inlinestaticconstexpr

◆ GetName()

template<typename ADataType_, typename AQDataType_, typename BDataType_, typename BQDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, typename QuantGroupSize_, bool TransposeC_, typename ComputeDataType_ = BDataType_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full>
CK_TILE_HOST const std::string ck_tile::GemmQuantPipelineProblemBase< ADataType_, AQDataType_, BDataType_, BQDataType_, CDataType_, BlockGemmShape_, Traits_, QuantGroupSize_, TransposeC_, ComputeDataType_, Scheduler_, HasHotLoop_, TailNum_ >::GetName ( )
inlinestaticnodiscard

Member Data Documentation

◆ DoubleSmemBuffer

template<typename ADataType_, typename AQDataType_, typename BDataType_, typename BQDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, typename QuantGroupSize_, bool TransposeC_, typename ComputeDataType_ = BDataType_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full>
bool ck_tile::GemmQuantPipelineProblemBase< ADataType_, AQDataType_, BDataType_, BQDataType_, CDataType_, BlockGemmShape_, Traits_, QuantGroupSize_, TransposeC_, ComputeDataType_, Scheduler_, HasHotLoop_, TailNum_ >::DoubleSmemBuffer = Traits::DoubleSmemBuffer
staticconstexpr

◆ HasHotLoop

template<typename ADataType_, typename AQDataType_, typename BDataType_, typename BQDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, typename QuantGroupSize_, bool TransposeC_, typename ComputeDataType_ = BDataType_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full>
auto ck_tile::GemmQuantPipelineProblemBase< ADataType_, AQDataType_, BDataType_, BQDataType_, CDataType_, BlockGemmShape_, Traits_, QuantGroupSize_, TransposeC_, ComputeDataType_, Scheduler_, HasHotLoop_, TailNum_ >::HasHotLoop = HasHotLoop_
staticconstexpr

◆ kBlockSize

template<typename ADataType_, typename AQDataType_, typename BDataType_, typename BQDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, typename QuantGroupSize_, bool TransposeC_, typename ComputeDataType_ = BDataType_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full>
index_t ck_tile::GemmPipelineProblemBase< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, ComputeDataType_, ck_tile::element_wise::PassThrough, ck_tile::element_wise::PassThrough, false, 1, 1 >::kBlockSize
staticconstexpr

◆ kPadK

template<typename ADataType_, typename AQDataType_, typename BDataType_, typename BQDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, typename QuantGroupSize_, bool TransposeC_, typename ComputeDataType_ = BDataType_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full>
bool ck_tile::GemmPipelineProblemBase< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, ComputeDataType_, ck_tile::element_wise::PassThrough, ck_tile::element_wise::PassThrough, false, 1, 1 >::kPadK
staticconstexpr

◆ kPadM

template<typename ADataType_, typename AQDataType_, typename BDataType_, typename BQDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, typename QuantGroupSize_, bool TransposeC_, typename ComputeDataType_ = BDataType_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full>
bool ck_tile::GemmPipelineProblemBase< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, ComputeDataType_, ck_tile::element_wise::PassThrough, ck_tile::element_wise::PassThrough, false, 1, 1 >::kPadM
staticconstexpr

◆ kPadN

template<typename ADataType_, typename AQDataType_, typename BDataType_, typename BQDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, typename QuantGroupSize_, bool TransposeC_, typename ComputeDataType_ = BDataType_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full>
bool ck_tile::GemmPipelineProblemBase< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, ComputeDataType_, ck_tile::element_wise::PassThrough, ck_tile::element_wise::PassThrough, false, 1, 1 >::kPadN
staticconstexpr

◆ PreshuffleB

template<typename ADataType_, typename AQDataType_, typename BDataType_, typename BQDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, typename QuantGroupSize_, bool TransposeC_, typename ComputeDataType_ = BDataType_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full>
bool ck_tile::GemmQuantPipelineProblemBase< ADataType_, AQDataType_, BDataType_, BQDataType_, CDataType_, BlockGemmShape_, Traits_, QuantGroupSize_, TransposeC_, ComputeDataType_, Scheduler_, HasHotLoop_, TailNum_ >::PreshuffleB = Traits::PreshuffleB
staticconstexpr

◆ Scheduler

template<typename ADataType_, typename AQDataType_, typename BDataType_, typename BQDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, typename QuantGroupSize_, bool TransposeC_, typename ComputeDataType_ = BDataType_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full>
auto ck_tile::GemmQuantPipelineProblemBase< ADataType_, AQDataType_, BDataType_, BQDataType_, CDataType_, BlockGemmShape_, Traits_, QuantGroupSize_, TransposeC_, ComputeDataType_, Scheduler_, HasHotLoop_, TailNum_ >::Scheduler = Scheduler_
staticconstexpr

◆ TailNum

template<typename ADataType_, typename AQDataType_, typename BDataType_, typename BQDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, typename QuantGroupSize_, bool TransposeC_, typename ComputeDataType_ = BDataType_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full>
auto ck_tile::GemmQuantPipelineProblemBase< ADataType_, AQDataType_, BDataType_, BQDataType_, CDataType_, BlockGemmShape_, Traits_, QuantGroupSize_, TransposeC_, ComputeDataType_, Scheduler_, HasHotLoop_, TailNum_ >::TailNum = TailNum_
staticconstexpr

◆ TransposeC

template<typename ADataType_, typename AQDataType_, typename BDataType_, typename BQDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, typename QuantGroupSize_, bool TransposeC_, typename ComputeDataType_ = BDataType_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full>
bool ck_tile::GemmQuantPipelineProblemBase< ADataType_, AQDataType_, BDataType_, BQDataType_, CDataType_, BlockGemmShape_, Traits_, QuantGroupSize_, TransposeC_, ComputeDataType_, Scheduler_, HasHotLoop_, TailNum_ >::TransposeC = TransposeC_
staticconstexpr

◆ VectorLoadSize

template<typename ADataType_, typename AQDataType_, typename BDataType_, typename BQDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, typename QuantGroupSize_, bool TransposeC_, typename ComputeDataType_ = BDataType_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full>
index_t ck_tile::GemmPipelineProblemBase< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, ComputeDataType_, ck_tile::element_wise::PassThrough, ck_tile::element_wise::PassThrough, false, 1, 1 >::VectorLoadSize
staticconstexpr

◆ VectorSizeAQ

template<typename ADataType_, typename AQDataType_, typename BDataType_, typename BQDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, typename QuantGroupSize_, bool TransposeC_, typename ComputeDataType_ = BDataType_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full>
index_t ck_tile::GemmQuantPipelineProblemBase< ADataType_, AQDataType_, BDataType_, BQDataType_, CDataType_, BlockGemmShape_, Traits_, QuantGroupSize_, TransposeC_, ComputeDataType_, Scheduler_, HasHotLoop_, TailNum_ >::VectorSizeAQ
staticconstexpr
Initial value:
= []() {
static_assert(std::is_same_v<ALayout, tensor_layout::gemm::RowMajor>);
return kPadK ? 1 : GetAlignmentAQ();
}()
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentAQ()
Definition gemm_quant_pipeline_problem.hpp:90
static constexpr bool kPadK
Definition gemm_pipeline_problem.hpp:80

◆ VectorSizeBQ

template<typename ADataType_, typename AQDataType_, typename BDataType_, typename BQDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, typename QuantGroupSize_, bool TransposeC_, typename ComputeDataType_ = BDataType_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full>
index_t ck_tile::GemmQuantPipelineProblemBase< ADataType_, AQDataType_, BDataType_, BQDataType_, CDataType_, BlockGemmShape_, Traits_, QuantGroupSize_, TransposeC_, ComputeDataType_, Scheduler_, HasHotLoop_, TailNum_ >::VectorSizeBQ = []() { return kPadK ? 1 : GetAlignmentBQ(); }()
staticconstexpr

The documentation for this struct was generated from the following file: