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 constexpr CK_TILE_HOST_DEVICE auto GetAlignmentA ()
 
static constexpr CK_TILE_HOST_DEVICE auto GetAlignmentB ()
 
static constexpr CK_TILE_HOST_DEVICE 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_>
static constexpr CK_TILE_HOST_DEVICE 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_>
static constexpr CK_TILE_HOST_DEVICE 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_>
static constexpr CK_TILE_HOST_DEVICE 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_>
static CK_TILE_HOST const std::string ck_tile::FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::GetName ( )
inlinestatic

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_>
constexpr 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_>
constexpr 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_>
constexpr 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_>
constexpr 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_>
constexpr 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_>
constexpr 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_>
constexpr 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_>
constexpr 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_>
constexpr 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_>
constexpr 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_>
constexpr 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_>
constexpr 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_>
constexpr 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 constexpr CK_TILE_HOST_DEVICE 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_>
constexpr 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 constexpr CK_TILE_HOST_DEVICE 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_>
constexpr 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 constexpr CK_TILE_HOST_DEVICE auto GetAlignmentC()
Definition: gemm_pipeline_problem.hpp:401
The documentation for this struct was generated from the following file:
  • /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_problem.hpp