UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ > Struct Template Reference

UniversalGemmPipelineProblem&lt; AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ &gt; Struct Template Reference#

Composable Kernel: ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ > Struct Template Reference
Public Types | Static Public Member Functions | Static Public Attributes | List of all members
ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ > Struct Template Reference

#include <gemm_pipeline_problem.hpp>

Public Types

using Traits = remove_cvref_t<Traits_>
using AsDataType = remove_cvref_t<AsDataType_>
using BsDataType = remove_cvref_t<BsDataType_>
using CDataType = remove_cvref_t<EDataType_>
using AElementWise = remove_cvref_t<AElementWise_>
using BElementWise = remove_cvref_t<BElementWise_>
using BlockGemmShape = remove_cvref_t<BlockGemmShape_>
using AsLayout = remove_cvref_t<typename Traits::AsLayout>
using BsLayout = remove_cvref_t<typename Traits::BsLayout>
using CLayout = remove_cvref_t<typename Traits::CLayout>
using ComputeDataTypeTuple
using AsLayoutTuple
using BsLayoutTuple
using AsDataTypeTuple
using BsDataTypeTuple
using ComputeDataType = remove_cvref_t<std::tuple_element_t<number<0>{}, ComputeDataTypeTuple>>
using ADataType = remove_cvref_t<std::tuple_element_t<number<0>{}, AsDataTypeTuple>>
using ALayout = remove_cvref_t<std::tuple_element_t<number<0>{}, AsLayoutTuple>>
using BDataType = remove_cvref_t<std::tuple_element_t<number<0>{}, BsDataTypeTuple>>
using BLayout = remove_cvref_t<std::tuple_element_t<number<0>{}, BsLayoutTuple>>

Static Public Member Functions

static CK_TILE_HOST const std::string GetName ()

Static Public Attributes

static constexpr bool FixedVectorSize = FixedVectorSize_
static constexpr bool ComputeDataTypeIsTuple = is_detected<is_tuple, ComputeDataType_>::value
static constexpr bool ADataTypeIsTuple = is_detected<is_tuple, AsDataType>::value
static constexpr bool BDataTypeIsTuple = is_detected<is_tuple, BsDataType>::value
static constexpr bool ALayoutIsTuple = is_detected<is_tuple, AsLayout>::value
static constexpr bool BLayoutIsTuple = is_detected<is_tuple, BsLayout>::value
static constexpr bool TransposeC = Traits::TransposeC
static constexpr index_t NumWaveGroups = Traits::NumWaveGroups
static constexpr bool UseStructuredSparsity = Traits::UseStructuredSparsity
static constexpr index_t kBlockSize = BlockGemmShape::NumWarps * get_warp_size()
static constexpr bool kPadM = Traits::kPadM
static constexpr bool kPadN = Traits::kPadN
static constexpr bool kPadK = Traits::kPadK
static constexpr bool DoubleSmemBuffer = Traits::DoubleSmemBuffer
static constexpr auto Scheduler = Scheduler_
static constexpr bool Preshuffle = Traits::Preshuffle
static constexpr index_t VectorSizeA = VectorSizeA_
static constexpr index_t VectorSizeB = VectorSizeB_
static constexpr auto HasHotLoop = HasHotLoop_
static constexpr auto TailNum = TailNum_
static constexpr index_t VectorLoadSize = Traits::_VectorSize

Member Typedef Documentation

◆ ADataType

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::ADataType = remove_cvref_t<std::tuple_element_t<number<0>{}, AsDataTypeTuple>>

◆ AElementWise

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::AElementWise = remove_cvref_t<AElementWise_>

◆ ALayout

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::ALayout = remove_cvref_t<std::tuple_element_t<number<0>{}, AsLayoutTuple>>

◆ AsDataType

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::AsDataType = remove_cvref_t<AsDataType_>

◆ AsDataTypeTuple

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::AsDataTypeTuple
Initial value:
std::conditional_t<ADataTypeIsTuple,
remove_cvref_t<AsDataType>,
remove_cvref_t<tuple<AsDataType>>>
ck_tile::remove_cvref_t
remove_cv_t< std::remove_reference_t< T > > remove_cvref_t
Definition type_traits.hpp:21
ck_tile::CShuffleEpilogue::ADataTypeIsTuple
static constexpr bool ADataTypeIsTuple
Definition cshuffle_epilogue.hpp:81

◆ AsLayout

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::AsLayout = remove_cvref_t<typename Traits::AsLayout>

◆ AsLayoutTuple

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::AsLayoutTuple
Initial value:
std::
conditional_t<ALayoutIsTuple, remove_cvref_t<AsLayout>, remove_cvref_t<tuple<AsLayout>>>

◆ BDataType

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::BDataType = remove_cvref_t<std::tuple_element_t<number<0>{}, BsDataTypeTuple>>

◆ BElementWise

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::BElementWise = remove_cvref_t<BElementWise_>

◆ BLayout

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::BLayout = remove_cvref_t<std::tuple_element_t<number<0>{}, BsLayoutTuple>>

◆ BlockGemmShape

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::BlockGemmShape = remove_cvref_t<BlockGemmShape_>

◆ BsDataType

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::BsDataType = remove_cvref_t<BsDataType_>

◆ BsDataTypeTuple

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::BsDataTypeTuple
Initial value:
std::conditional_t<BDataTypeIsTuple,
remove_cvref_t<BsDataType>,
remove_cvref_t<tuple<BsDataType>>>
ck_tile::CShuffleEpilogue::BDataTypeIsTuple
static constexpr bool BDataTypeIsTuple
Definition cshuffle_epilogue.hpp:82

◆ BsLayout

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::BsLayout = remove_cvref_t<typename Traits::BsLayout>

◆ BsLayoutTuple

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::BsLayoutTuple
Initial value:
std::
conditional_t<BLayoutIsTuple, remove_cvref_t<BsLayout>, remove_cvref_t<tuple<BsLayout>>>

◆ CDataType

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::CDataType = remove_cvref_t<EDataType_>

◆ CLayout

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::CLayout = remove_cvref_t<typename Traits::CLayout>

◆ ComputeDataType

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::ComputeDataType = remove_cvref_t<std::tuple_element_t<number<0>{}, ComputeDataTypeTuple>>

◆ ComputeDataTypeTuple

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::ComputeDataTypeTuple
Initial value:
std::conditional_t<ComputeDataTypeIsTuple,
remove_cvref_t<ComputeDataType_>,
remove_cvref_t<tuple<ComputeDataType_>>>
ck_tile::GemmPipelineProblemBase::ComputeDataTypeIsTuple
static constexpr bool ComputeDataTypeIsTuple
Definition gemm_pipeline_problem.hpp:43

◆ Traits

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::Traits = remove_cvref_t<Traits_>

Member Function Documentation

◆ GetName()

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
CK_TILE_HOST const std::string ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::GetName ( )
inlinestaticnodiscard

Member Data Documentation

◆ ADataTypeIsTuple

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
bool ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::ADataTypeIsTuple = is_detected<is_tuple, AsDataType>::value
staticconstexpr

◆ ALayoutIsTuple

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
bool ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::ALayoutIsTuple = is_detected<is_tuple, AsLayout>::value
staticconstexpr

◆ BDataTypeIsTuple

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
bool ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::BDataTypeIsTuple = is_detected<is_tuple, BsDataType>::value
staticconstexpr

◆ BLayoutIsTuple

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
bool ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::BLayoutIsTuple = is_detected<is_tuple, BsLayout>::value
staticconstexpr

◆ ComputeDataTypeIsTuple

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
bool ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::ComputeDataTypeIsTuple = is_detected<is_tuple, ComputeDataType_>::value
staticconstexpr

◆ DoubleSmemBuffer

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
bool ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::DoubleSmemBuffer = Traits::DoubleSmemBuffer
staticconstexpr

◆ FixedVectorSize

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
bool ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::FixedVectorSize = FixedVectorSize_
staticconstexpr

◆ HasHotLoop

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
auto ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::HasHotLoop = HasHotLoop_
staticconstexpr

◆ kBlockSize

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
index_t ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::kBlockSize = BlockGemmShape::NumWarps * get_warp_size()
staticconstexpr

◆ kPadK

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
bool ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::kPadK = Traits::kPadK
staticconstexpr

◆ kPadM

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
bool ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::kPadM = Traits::kPadM
staticconstexpr

◆ kPadN

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
bool ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::kPadN = Traits::kPadN
staticconstexpr

◆ NumWaveGroups

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
index_t ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::NumWaveGroups = Traits::NumWaveGroups
staticconstexpr

◆ Preshuffle

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
bool ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::Preshuffle = Traits::Preshuffle
staticconstexpr

◆ Scheduler

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
auto ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::Scheduler = Scheduler_
staticconstexpr

◆ TailNum

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
auto ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::TailNum = TailNum_
staticconstexpr

◆ TransposeC

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
bool ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::TransposeC = Traits::TransposeC
staticconstexpr

◆ UseStructuredSparsity

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
bool ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::UseStructuredSparsity = Traits::UseStructuredSparsity
staticconstexpr

◆ VectorLoadSize

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
index_t ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::VectorLoadSize = Traits::_VectorSize
staticconstexpr

◆ VectorSizeA

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
index_t ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::VectorSizeA = VectorSizeA_
staticconstexpr

◆ VectorSizeB

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, typename ComputeDataType_ = AsDataType_, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
index_t ck_tile::UniversalGemmPipelineProblem< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, AElementWise_, BElementWise_, ComputeDataType_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::VectorSizeB = VectorSizeB_
staticconstexpr
The documentation for this struct was generated from the following file: