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

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

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

#include <mixed_prec_flatmm_pipeline_agmem_bgmem_creg_v1.hpp>

Inheritance diagram for ck_tile::F16xMXF4FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >:
ck_tile::FlatmmPipelineProblem< ADataType_, ADataType_, CDataType_, BlockGemmShape_, Traits_, GemmPipelineScheduler::Intrawave, true, TailNumber::Full, ADataType_ >

Public Types

using BlockGemmShape = BlockGemmShape_
using QuantType = BDataType_
Public Types inherited from ck_tile::FlatmmPipelineProblem< ADataType_, ADataType_, CDataType_, BlockGemmShape_, Traits_, GemmPipelineScheduler::Intrawave, true, TailNumber::Full, ADataType_ >
using Traits
using ADataType
using BDataType
using CDataType
using ComputeDataType
using BlockGemmShape
using ALayout
using BLayout
using CLayout

Static Public Attributes

static constexpr index_t flatNPerWarp = BlockGemmShape::flatNPerWarp
static constexpr int MXF4ScaleGranularityK = 32
static constexpr int ContinuousKPerThread = 32
static constexpr int ContinuousScaleNPerThread = 2
static constexpr int ContinuousScaleKPerThread = 2
static constexpr index_t flatKPerWarp = 64 * ContinuousKPerThread
Static Public Attributes inherited from ck_tile::FlatmmPipelineProblem< ADataType_, ADataType_, CDataType_, BlockGemmShape_, Traits_, GemmPipelineScheduler::Intrawave, true, TailNumber::Full, ADataType_ >
static constexpr bool TransposeC
static constexpr index_t NumWaveGroups
static constexpr bool UseStructuredSparsity
static constexpr index_t kBlockSize
static constexpr bool kPadM
static constexpr bool kPadN
static constexpr bool kPadK
static constexpr bool DoubleSmemBuffer
static constexpr auto Scheduler
static constexpr index_t VectorLoadSize
static constexpr auto HasHotLoop
static constexpr auto TailNum
static constexpr index_t VectorSizeA
static constexpr index_t VectorSizeB
static constexpr index_t VectorSizeC

Additional Inherited Members

Static Public Member Functions inherited from ck_tile::FlatmmPipelineProblem< ADataType_, ADataType_, CDataType_, BlockGemmShape_, Traits_, GemmPipelineScheduler::Intrawave, true, TailNumber::Full, ADataType_ >
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 ()

Member Typedef Documentation

◆ 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::F16xMXF4FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::BlockGemmShape = BlockGemmShape_

◆ QuantType

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::F16xMXF4FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::QuantType = BDataType_

Member Data Documentation

◆ ContinuousKPerThread

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_>
int ck_tile::F16xMXF4FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::ContinuousKPerThread = 32
staticconstexpr

◆ ContinuousScaleKPerThread

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_>
int ck_tile::F16xMXF4FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::ContinuousScaleKPerThread = 2
staticconstexpr

◆ ContinuousScaleNPerThread

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_>
int ck_tile::F16xMXF4FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::ContinuousScaleNPerThread = 2
staticconstexpr

◆ flatKPerWarp

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::F16xMXF4FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::flatKPerWarp = 64 * ContinuousKPerThread
staticconstexpr

◆ flatNPerWarp

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::F16xMXF4FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::flatNPerWarp = BlockGemmShape::flatNPerWarp
staticconstexpr

◆ MXF4ScaleGranularityK

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_>
int ck_tile::F16xMXF4FlatmmPipelineProblem< ADataType_, BDataType_, CDataType_, BlockGemmShape_, Traits_, Scheduler_, HasHotLoop_, TailNum_, ComputeDataType_ >::MXF4ScaleGranularityK = 32
staticconstexpr
The documentation for this struct was generated from the following file: