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_ >:
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:
GemmPipelineProblemBase<ADataType_,
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>
|
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>
|
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>
|
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>
|
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>
|
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>
|
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>
|
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>
|
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>
|
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>
|
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>
|
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>
|
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>
|
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>
|
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>
|
staticconstexpr |
Initial value:
= []() {
static_assert(std::is_same_v<ALayout, tensor_layout::gemm::RowMajor>);
}()
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>
|
staticconstexpr |
The documentation for this struct was generated from the following file: