GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ > Struct Template Reference

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

Composable Kernel: ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ > Struct Template Reference
Public Types | Static Public Member Functions | Static Public Attributes | List of all members
ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, 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 BlockGemmShape = remove_cvref_t<BlockGemmShape_>
using AElementWise = remove_cvref_t<AElementWise_>
using BElementWise = remove_cvref_t<BElementWise_>
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 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 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 = GemmPipelineScheduler::Default
static constexpr index_t VectorLoadSize = Traits::_VectorSize
static constexpr bool Preshuffle = false
static constexpr index_t VectorSizeA
static constexpr index_t VectorSizeB
static constexpr index_t VectorSizeC

Member Typedef Documentation

◆ ADataType

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, 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_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::AElementWise = remove_cvref_t<AElementWise_>

◆ ALayout

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, 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_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::AsDataType = remove_cvref_t<AsDataType_>

◆ AsDataTypeTuple

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, 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_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::AsLayout = remove_cvref_t<typename Traits::AsLayout>

◆ AsLayoutTuple

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, 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_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, 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_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::BElementWise = remove_cvref_t<BElementWise_>

◆ BLayout

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, 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_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::BlockGemmShape = remove_cvref_t<BlockGemmShape_>

◆ BsDataType

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::BsDataType = remove_cvref_t<BsDataType_>

◆ BsDataTypeTuple

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, 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_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::BsLayout = remove_cvref_t<typename Traits::BsLayout>

◆ BsLayoutTuple

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, 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_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::CDataType = remove_cvref_t<EDataType_>

◆ CLayout

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::CLayout = remove_cvref_t<typename Traits::CLayout>

◆ ComputeDataType

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, 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_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, 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_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
using ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::Traits = remove_cvref_t<Traits_>

Member Function Documentation

◆ GetAlignmentA()

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::GetAlignmentA ( )
inlinestaticconstexpr

◆ GetAlignmentB()

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::GetAlignmentB ( )
inlinestaticconstexpr

◆ GetAlignmentC()

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::GetAlignmentC ( )
inlinestaticconstexpr

◆ GetName()

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
CK_TILE_HOST const std::string ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::GetName ( )
inlinestaticnodiscard

Member Data Documentation

◆ ADataTypeIsTuple

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
bool ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::ADataTypeIsTuple = is_detected<is_tuple, AsDataType>::value
staticconstexpr

◆ ALayoutIsTuple

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
bool ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::ALayoutIsTuple = is_detected<is_tuple, AsLayout>::value
staticconstexpr

◆ BDataTypeIsTuple

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
bool ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::BDataTypeIsTuple = is_detected<is_tuple, BsDataType>::value
staticconstexpr

◆ BLayoutIsTuple

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
bool ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::BLayoutIsTuple = is_detected<is_tuple, BsLayout>::value
staticconstexpr

◆ ComputeDataTypeIsTuple

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
bool ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::ComputeDataTypeIsTuple = is_detected<is_tuple, ComputeDataType_>::value
staticconstexpr

◆ DoubleSmemBuffer

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
bool ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::DoubleSmemBuffer = Traits::DoubleSmemBuffer
staticconstexpr

◆ FixedVectorSize

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
bool ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::FixedVectorSize = FixedVectorSize_
staticconstexpr

◆ kBlockSize

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
index_t ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::kBlockSize = BlockGemmShape::NumWarps * get_warp_size()
staticconstexpr

◆ kPadK

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
bool ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::kPadK = Traits::kPadK
staticconstexpr

◆ kPadM

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
bool ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::kPadM = Traits::kPadM
staticconstexpr

◆ kPadN

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
bool ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::kPadN = Traits::kPadN
staticconstexpr

◆ NumWaveGroups

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
index_t ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::NumWaveGroups = Traits::NumWaveGroups
staticconstexpr

◆ Preshuffle

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
bool ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::Preshuffle = false
staticconstexpr

◆ Scheduler

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
auto ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::Scheduler = GemmPipelineScheduler::Default
staticconstexpr

◆ TransposeC

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
bool ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::TransposeC = Traits::TransposeC
staticconstexpr

◆ UseStructuredSparsity

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
bool ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::UseStructuredSparsity = Traits::UseStructuredSparsity
staticconstexpr

◆ VectorLoadSize

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
index_t ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::VectorLoadSize = Traits::_VectorSize
staticconstexpr

◆ VectorSizeA

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
index_t ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::VectorSizeA
staticconstexpr
Initial value:
= []() {
if constexpr(FixedVectorSize)
{
return VectorSizeA_;
}
else if constexpr(std::is_same_v<AsLayout, tensor_layout::gemm::RowMajor>)
{
return kPadK ? 1 : GetAlignmentA();
}
else
{
return kPadM ? 1 : GetAlignmentA();
}
}()
ck_tile::GemmPipelineProblemBase::FixedVectorSize
static constexpr bool FixedVectorSize
Definition gemm_pipeline_problem.hpp:32
ck_tile::GemmPipelineProblemBase::kPadK
static constexpr bool kPadK
Definition gemm_pipeline_problem.hpp:80
ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::kPadM
static constexpr bool kPadM
Definition gemm_pipeline_problem.hpp:78
ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::GetAlignmentA
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentA()
Definition gemm_pipeline_problem.hpp:99

◆ VectorSizeB

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
index_t ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::VectorSizeB
staticconstexpr
Initial value:
= []() {
if constexpr(FixedVectorSize)
{
return VectorSizeB_;
}
else if constexpr(std::is_same_v<BsLayout, tensor_layout::gemm::ColumnMajor>)
{
return kPadN ? 1 : GetAlignmentB();
}
else
{
return kPadK ? 1 : GetAlignmentB();
}
}()
ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::GetAlignmentB
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentB()
Definition gemm_pipeline_problem.hpp:117
ck_tile::GemmPipelineProblemBase::kPadN
static constexpr bool kPadN
Definition gemm_pipeline_problem.hpp:79

◆ VectorSizeC

template<typename AsDataType_, typename BsDataType_, typename EDataType_, typename BlockGemmShape_, typename Traits_, typename ComputeDataType_ = AsDataType_, typename AElementWise_ = ck_tile::element_wise::PassThrough, typename BElementWise_ = ck_tile::element_wise::PassThrough, bool FixedVectorSize_ = false, index_t VectorSizeA_ = 1, index_t VectorSizeB_ = 1>
index_t ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::VectorSizeC
staticconstexpr
Initial value:
= []() {
if constexpr(std::is_same_v<CLayout, tensor_layout::gemm::RowMajor>)
{
return kPadN ? 1 : GetAlignmentC();
}
else
{
return kPadM ? 1 : GetAlignmentC();
}
}()
ck_tile::GemmPipelineProblemBase< AsDataType_, BsDataType_, EDataType_, BlockGemmShape_, Traits_, ComputeDataType_, AElementWise_, BElementWise_, FixedVectorSize_, VectorSizeA_, VectorSizeB_ >::GetAlignmentC
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentC()
Definition gemm_pipeline_problem.hpp:135
The documentation for this struct was generated from the following file: