FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ > Struct Template Reference

FlatmmPipelineProblem&lt; ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ &gt; Struct Template Reference#

Composable Kernel: ck_tile::FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ > Struct Template Reference
Public Types | Static Public Member Functions | Static Public Attributes | List of all members
ck_tile::FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ > Struct Template Reference

#include <gemm_pipeline_problem.hpp>

Public Types

using Traits = remove_cvref_t<Traits_>
using ADataType = remove_cvref_t<ADataType_>
using BDataType = remove_cvref_t<BDataType_>
using CDataType = remove_cvref_t<CDataType_>
using ComputeDataType = remove_cvref_t<ComputeDataType_>
using BlockGemmShape = remove_cvref_t<BlockGemmShape_>
using ALayout = remove_cvref_t<typename Traits::AsLayout>
using BLayout = remove_cvref_t<typename Traits::BsLayout>
using CLayout = remove_cvref_t<typename Traits::CLayout>

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 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 auto HasHotLoop = HasHotLoop_
static constexpr auto TailNum = TailNum_
static constexpr index_t VectorSizeA
static constexpr index_t VectorSizeB
static constexpr index_t VectorSizeC

Member Typedef Documentation

◆ ADataType

template<typename ADataType_, typename BDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename ComputeDataType_ = ADataType_>
using ck_tile::FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::ADataType = remove_cvref_t<ADataType_>

◆ ALayout

template<typename ADataType_, typename BDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename ComputeDataType_ = ADataType_>
using ck_tile::FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::ALayout = remove_cvref_t<typename Traits::AsLayout>

◆ BDataType

template<typename ADataType_, typename BDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename ComputeDataType_ = ADataType_>
using ck_tile::FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::BDataType = remove_cvref_t<BDataType_>

◆ BLayout

template<typename ADataType_, typename BDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename ComputeDataType_ = ADataType_>
using ck_tile::FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::BLayout = remove_cvref_t<typename Traits::BsLayout>

◆ BlockGemmShape

template<typename ADataType_, typename BDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename ComputeDataType_ = ADataType_>
using ck_tile::FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::BlockGemmShape = remove_cvref_t<BlockGemmShape_>

◆ CDataType

template<typename ADataType_, typename BDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename ComputeDataType_ = ADataType_>
using ck_tile::FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::CDataType = remove_cvref_t<CDataType_>

◆ CLayout

template<typename ADataType_, typename BDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename ComputeDataType_ = ADataType_>
using ck_tile::FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::CLayout = remove_cvref_t<typename Traits::CLayout>

◆ ComputeDataType

template<typename ADataType_, typename BDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename ComputeDataType_ = ADataType_>
using ck_tile::FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::ComputeDataType = remove_cvref_t<ComputeDataType_>

◆ Traits

template<typename ADataType_, typename BDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename ComputeDataType_ = ADataType_>
using ck_tile::FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::Traits = remove_cvref_t<Traits_>

Member Function Documentation

◆ GetAlignmentA()

template<typename ADataType_, typename BDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename ComputeDataType_ = ADataType_>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::GetAlignmentA ( )
inlinestaticconstexpr

◆ GetAlignmentB()

template<typename ADataType_, typename BDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename ComputeDataType_ = ADataType_>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::GetAlignmentB ( )
inlinestaticconstexpr

◆ GetAlignmentC()

template<typename ADataType_, typename BDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename ComputeDataType_ = ADataType_>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::GetAlignmentC ( )
inlinestaticconstexpr

◆ GetName()

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

Member Data Documentation

◆ DoubleSmemBuffer

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

◆ HasHotLoop

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

◆ kBlockSize

template<typename ADataType_, typename BDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename ComputeDataType_ = ADataType_>
index_t ck_tile::FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::kBlockSize = BlockGemmShape::NumWarps * get_warp_size()
staticconstexpr

◆ kPadK

template<typename ADataType_, typename BDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename ComputeDataType_ = ADataType_>
bool ck_tile::FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::kPadK = Traits::kPadK
staticconstexpr

◆ kPadM

template<typename ADataType_, typename BDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename ComputeDataType_ = ADataType_>
bool ck_tile::FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::kPadM = Traits::kPadM
staticconstexpr

◆ kPadN

template<typename ADataType_, typename BDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename ComputeDataType_ = ADataType_>
bool ck_tile::FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::kPadN = Traits::kPadN
staticconstexpr

◆ NumWaveGroups

template<typename ADataType_, typename BDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename ComputeDataType_ = ADataType_>
index_t ck_tile::FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::NumWaveGroups = Traits::NumWaveGroups
staticconstexpr

◆ Scheduler

template<typename ADataType_, typename BDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename ComputeDataType_ = ADataType_>
auto ck_tile::FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::Scheduler = GemmPipelineScheduler::Default
staticconstexpr

◆ TailNum

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

◆ TransposeC

template<typename ADataType_, typename BDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename ComputeDataType_ = ADataType_>
bool ck_tile::FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::TransposeC = Traits::TransposeC
staticconstexpr

◆ UseStructuredSparsity

template<typename ADataType_, typename BDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename ComputeDataType_ = ADataType_>
bool ck_tile::FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::UseStructuredSparsity = Traits::UseStructuredSparsity
staticconstexpr

◆ VectorLoadSize

template<typename ADataType_, typename BDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename ComputeDataType_ = ADataType_>
index_t ck_tile::FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::VectorLoadSize = Traits::_VectorSize
staticconstexpr

◆ VectorSizeA

template<typename ADataType_, typename BDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename ComputeDataType_ = ADataType_>
index_t ck_tile::FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::VectorSizeA
staticconstexpr
Initial value:
= []() {
if constexpr(std::is_same_v<ALayout, tensor_layout::gemm::RowMajor>)
{
return kPadK ? 1 : GetAlignmentA();
}
else
{
return kPadM ? 1 : GetAlignmentA();
}
}()
ck_tile::FlatmmPipelineProblem::GetAlignmentA
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentA()
Definition gemm_pipeline_problem.hpp:365
ck_tile::FlatmmPipelineProblem::kPadM
static constexpr bool kPadM
Definition gemm_pipeline_problem.hpp:343
ck_tile::FlatmmPipelineProblem::kPadK
static constexpr bool kPadK
Definition gemm_pipeline_problem.hpp:345

◆ VectorSizeB

template<typename ADataType_, typename BDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename ComputeDataType_ = ADataType_>
index_t ck_tile::FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::VectorSizeB
staticconstexpr
Initial value:
= []() {
if constexpr(std::is_same_v<BLayout, tensor_layout::gemm::ColumnMajor>)
{
return kPadN ? 1 : GetAlignmentB();
}
else
{
return kPadK ? 1 : GetAlignmentB();
}
}()
ck_tile::FlatmmPipelineProblem::kPadN
static constexpr bool kPadN
Definition gemm_pipeline_problem.hpp:344
ck_tile::FlatmmPipelineProblem::GetAlignmentB
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentB()
Definition gemm_pipeline_problem.hpp:383

◆ VectorSizeC

template<typename ADataType_, typename BDataType_, typename CDataType_, typename BlockGemmShape_, typename Traits_, GemmPipelineScheduler Scheduler_ = GemmPipelineScheduler::Intrawave, bool HasHotLoop_ = true, TailNumber TailNum_ = TailNumber::Full, typename ComputeDataType_ = ADataType_>
index_t ck_tile::FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::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::FlatmmPipelineProblem::GetAlignmentC
static CK_TILE_HOST_DEVICE constexpr auto GetAlignmentC()
Definition gemm_pipeline_problem.hpp:401
ck_tile::GemmPipelineProblemBase::kPadN
static constexpr bool kPadN
Definition gemm_pipeline_problem.hpp:79
The documentation for this struct was generated from the following file: