BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ > Struct Template Reference

BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ > Struct Template Reference#

Composable Kernel: ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ > Struct Template Reference

#include <block_fmha_pipeline_qr_ks_vs_async_trload.hpp>

Public Types
using	Problem = remove_cvref_t<Problem_>
using	Policy = remove_cvref_t<Policy_>
using	QDataType = remove_cvref_t<typename Problem::QDataType>
using	KDataType = remove_cvref_t<typename Problem::KDataType>
using	VDataType = remove_cvref_t<typename Problem::VDataType>
using	SaccDataType = remove_cvref_t<typename Problem::SaccDataType>
using	SMPLComputeDataType = remove_cvref_t<typename Problem::SMPLComputeDataType>
using	BiasDataType = remove_cvref_t<typename Problem::BiasDataType>
using	RandValOutputDataType = remove_cvref_t<typename Problem::RandValOutputDataType>
using	LSEDataType = remove_cvref_t<typename Problem::LSEDataType>
using	PDataType = remove_cvref_t<typename Problem::PDataType>
using	OaccDataType = remove_cvref_t<typename Problem::OaccDataType>
using	ODataType = remove_cvref_t<typename Problem::ODataType>
using	AttentionVariant = remove_cvref_t<typename Problem::AttentionVariant>
using	FmhaMask = remove_cvref_t<typename Problem::FmhaMask>
using	BlockFmhaShape = remove_cvref_t<typename Problem::BlockFmhaShape>
using	VLayout = remove_cvref_t<typename BlockFmhaShape::VLayout>

Public Member Functions
template<typename QDramBlockWindowTmp, typename KDramBlockWindowTmp, typename VDramBlockWindowTmp, typename BiasDramBlockWindowTmp, typename LSEaccDramBlockWindowTmp, typename PositionEncoding>
CK_TILE_HOST_DEVICE auto	operator() (const QDramBlockWindowTmp &q_dram_block_window_tmp, const KDramBlockWindowTmp &k_dram_block_window_tmp, const VDramBlockWindowTmp &v_dram_block_window_tmp, const BiasDramBlockWindowTmp &bias_dram_block_window_tmp, LSEaccDramBlockWindowTmp &lse_acc_dram_window_tmp, FmhaMask mask, PositionEncoding position_encoding, float scale_s, void *smem_ptr) const
template<typename QDramBlockWindowTmp, typename KDramBlockWindowTmp, typename VDramBlockWindowTmp, typename BiasDramBlockWindowTmp, typename LSEaccDramBlockWindowTmp, typename PositionEncoding>
CK_TILE_HOST_DEVICE auto	operator() (const QDramBlockWindowTmp &__restrict__ q_dram_block_window_tmp, const KDramBlockWindowTmp &__restrict__ k_dram_block_window_tmp, const VDramBlockWindowTmp &__restrict__ v_dram_block_window_tmp, const BiasDramBlockWindowTmp &__restrict__ bias_dram_block_window_tmp, LSEaccDramBlockWindowTmp &__restrict__ lse_acc_dram_window_tmp, FmhaMask mask, PositionEncoding position_encoding, float scale_s, void __restrict__ smem_ptrk0, void __restrict__ smem_ptrk1, void __restrict__ smem_ptrv0, void __restrict__ smem_ptrv1) const

Static Public Member Functions
static CK_TILE_HOST_DEVICE constexpr ck_tile::index_t	GetSmemSize ()

Static Public Attributes
static constexpr auto	I0 = number<0>{}
static constexpr auto	I1 = number<1>{}
static constexpr bool	kQLoadOnce = true
static constexpr bool	kKLoadOnce = BlockFmhaShape::kM0 >= 64
static constexpr index_t	kBlockSize = Problem::kBlockSize
static constexpr index_t	kM0 = BlockFmhaShape::kM0
static constexpr index_t	kN0 = BlockFmhaShape::kN0
static constexpr index_t	kK0 = BlockFmhaShape::kK0
static constexpr index_t	kN1 = BlockFmhaShape::kN1
static constexpr index_t	kK1 = BlockFmhaShape::kK1
static constexpr index_t	kQKHeaddim = BlockFmhaShape::kQKHeaddim
static constexpr index_t	kSubQKHeaddim = BlockFmhaShape::kSubQKHeaddim
static constexpr index_t	kNWarp = BlockFmhaShape::Gemm0BlockWarps::at(I1)
static constexpr index_t	kNXdl = BlockFmhaShape::Gemm0WarpTile::at(I1)
static constexpr bool	kIsGroupMode = Problem::kIsGroupMode
static constexpr bool	kPadSeqLenQ = Problem::kPadSeqLenQ
static constexpr bool	kPadSeqLenK = Problem::kPadSeqLenK
static constexpr bool	kPadHeadDimQ
static constexpr bool	kPadHeadDimV
static constexpr bool	kHasLogitsSoftCap = Problem::kHasLogitsSoftCap
static constexpr bool	kHasDropout = Problem::kHasDropout
static constexpr auto	BiasEnum = Problem::BiasEnum
static constexpr bool	kStoreLSE = Problem::kStoreLSE
static constexpr bool	kHasUnevenSplits = true
static constexpr index_t	kAlignmentQ = Policy::template GetAlignmentQ<Problem>()
static constexpr index_t	kAlignmentK = Policy::template GetAlignmentK<Problem>()
static constexpr index_t	kAlignmentV
static constexpr index_t	kAlignmentOacc = Policy::template GetAlignmentO<Problem>()
static constexpr index_t	kAlignmentBias
static constexpr index_t	kBlockPerCu
static constexpr const char *	name = "qr_async_trload"

Member Typedef Documentation

◆ AttentionVariant

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

using ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::AttentionVariant = remove_cvref_t<typename Problem::AttentionVariant>

◆ BiasDataType

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

using ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::BiasDataType = remove_cvref_t<typename Problem::BiasDataType>

◆ BlockFmhaShape

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

using ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::BlockFmhaShape = remove_cvref_t<typename Problem::BlockFmhaShape>

◆ FmhaMask

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

using ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::FmhaMask = remove_cvref_t<typename Problem::FmhaMask>

◆ KDataType

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

using ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::KDataType = remove_cvref_t<typename Problem::KDataType>

◆ LSEDataType

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

using ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::LSEDataType = remove_cvref_t<typename Problem::LSEDataType>

◆ OaccDataType

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

using ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::OaccDataType = remove_cvref_t<typename Problem::OaccDataType>

◆ ODataType

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

using ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::ODataType = remove_cvref_t<typename Problem::ODataType>

◆ PDataType

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

using ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::PDataType = remove_cvref_t<typename Problem::PDataType>

◆ Policy

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

using ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::Policy = remove_cvref_t<Policy_>

◆ Problem

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

using ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::Problem = remove_cvref_t<Problem_>

◆ QDataType

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

using ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::QDataType = remove_cvref_t<typename Problem::QDataType>

◆ RandValOutputDataType

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

using ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::RandValOutputDataType = remove_cvref_t<typename Problem::RandValOutputDataType>

◆ SaccDataType

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

using ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::SaccDataType = remove_cvref_t<typename Problem::SaccDataType>

◆ SMPLComputeDataType

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

using ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::SMPLComputeDataType = remove_cvref_t<typename Problem::SMPLComputeDataType>

◆ VDataType

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

using ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::VDataType = remove_cvref_t<typename Problem::VDataType>

◆ VLayout

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

using ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::VLayout = remove_cvref_t<typename BlockFmhaShape::VLayout>

Member Function Documentation

◆ GetSmemSize()

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

CK_TILE_HOST_DEVICE constexpr ck_tile::index_t ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::GetSmemSize ( )

inlinestaticconstexpr

◆ operator()() [1/2]

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

template<typename QDramBlockWindowTmp, typename KDramBlockWindowTmp, typename VDramBlockWindowTmp, typename BiasDramBlockWindowTmp, typename LSEaccDramBlockWindowTmp, typename PositionEncoding>

CK_TILE_HOST_DEVICE auto ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::operator()	(	const QDramBlockWindowTmp &__restrict__	q_dram_block_window_tmp,
		const KDramBlockWindowTmp &__restrict__	k_dram_block_window_tmp,
		const VDramBlockWindowTmp &__restrict__	v_dram_block_window_tmp,
		const BiasDramBlockWindowTmp &__restrict__	bias_dram_block_window_tmp,
		LSEaccDramBlockWindowTmp &__restrict__	lse_acc_dram_window_tmp,
		FmhaMask	mask,
		PositionEncoding	position_encoding,
		float	scale_s,
		void *__restrict__	smem_ptrk0,
		void *__restrict__	smem_ptrk1,
		void *__restrict__	smem_ptrv0,
		void *__restrict__	smem_ptrv1 ) const

inline

NOTICE: bias might be materialized mask including -inf values, need consideration

◆ operator()() [2/2]

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

template<typename QDramBlockWindowTmp, typename KDramBlockWindowTmp, typename VDramBlockWindowTmp, typename BiasDramBlockWindowTmp, typename LSEaccDramBlockWindowTmp, typename PositionEncoding>

CK_TILE_HOST_DEVICE auto ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::operator()	(	const QDramBlockWindowTmp &	q_dram_block_window_tmp,
		const KDramBlockWindowTmp &	k_dram_block_window_tmp,
		const VDramBlockWindowTmp &	v_dram_block_window_tmp,
		const BiasDramBlockWindowTmp &	bias_dram_block_window_tmp,
		LSEaccDramBlockWindowTmp &	lse_acc_dram_window_tmp,
		FmhaMask	mask,
		PositionEncoding	position_encoding,
		float	scale_s,
		void *	smem_ptr ) const

inline

NOTICE: bias might be materialized mask including -inf values, need consideration

Member Data Documentation

◆ BiasEnum

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

auto ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::BiasEnum = Problem::BiasEnum

staticconstexpr

◆ I0

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

auto ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::I0 = number<0>{}

staticconstexpr

◆ I1

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

auto ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::I1 = number<1>{}

staticconstexpr

◆ kAlignmentBias

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

index_t ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::kAlignmentBias

staticconstexpr

Initial value:

=

kPadSeqLenK ? 1 : Policy::template GetAlignmentBias<Problem>()

ck_tile::BlockFmhaBatchPrefillPipelineQRKSVSAsync::kPadSeqLenK

static constexpr bool kPadSeqLenK

Definition block_fmha_batch_prefill_pipeline_qr_ks_vs_async.hpp:64

◆ kAlignmentK

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

index_t ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::kAlignmentK = Policy::template GetAlignmentK<Problem>()

staticconstexpr

◆ kAlignmentOacc

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

index_t ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::kAlignmentOacc = Policy::template GetAlignmentO<Problem>()

staticconstexpr

◆ kAlignmentQ

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

index_t ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::kAlignmentQ = Policy::template GetAlignmentQ<Problem>()

staticconstexpr

◆ kAlignmentV

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

index_t ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::kAlignmentV

staticconstexpr

Initial value:

= []() {
        if constexpr(std::is_same_v<VLayout, ck_tile::tensor_layout::gemm::RowMajor>)
            return Policy::template GetAlignmentV<Problem>();
        else
            return kPadSeqLenK ? 1 : Policy::template GetAlignmentV<Problem>();
    }()

◆ kBlockPerCu

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

index_t ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::kBlockPerCu

staticconstexpr

Initial value:

= []() {
        if constexpr(Problem::kBlockPerCu != -1)
            return Problem::kBlockPerCu;
        else
        {
            if constexpr(kQKHeaddim <= 32)
            {
                return 2;
            }
            else if constexpr(kQKHeaddim <= 64)
            {
                return 3;
            }
            else if constexpr(kQKHeaddim <= 128)
            {
                if constexpr(BiasEnum == BlockAttentionBiasEnum::ELEMENTWISE_BIAS || kM0 >= 256)
                    return 1;
                else
                    return 2;
            }
            else if constexpr(kQKHeaddim <= 256)
            {
                return 1;
            }
            else
            {
                return 1;
            }
        }
    }()

◆ kBlockSize

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

index_t ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::kBlockSize = Problem::kBlockSize

staticconstexpr

◆ kHasDropout

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

bool ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::kHasDropout = Problem::kHasDropout

staticconstexpr

◆ kHasLogitsSoftCap

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

bool ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::kHasLogitsSoftCap = Problem::kHasLogitsSoftCap

staticconstexpr

◆ kHasUnevenSplits

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

bool ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::kHasUnevenSplits = true

staticconstexpr

◆ kIsGroupMode

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

bool ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::kIsGroupMode = Problem::kIsGroupMode

staticconstexpr

◆ kK0

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

index_t ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::kK0 = BlockFmhaShape::kK0

staticconstexpr

◆ kK1

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

index_t ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::kK1 = BlockFmhaShape::kK1

staticconstexpr

◆ kKLoadOnce

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

bool ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::kKLoadOnce = BlockFmhaShape::kM0 >= 64

staticconstexpr

◆ kM0

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

index_t ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::kM0 = BlockFmhaShape::kM0

staticconstexpr

◆ kN0

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

index_t ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::kN0 = BlockFmhaShape::kN0

staticconstexpr

◆ kN1

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

index_t ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::kN1 = BlockFmhaShape::kN1

staticconstexpr

◆ kNWarp

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

index_t ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::kNWarp = BlockFmhaShape::Gemm0BlockWarps::at(I1)

staticconstexpr

◆ kNXdl

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

index_t ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::kNXdl = BlockFmhaShape::Gemm0WarpTile::at(I1)

staticconstexpr

◆ kPadHeadDimQ

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

bool ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::kPadHeadDimQ

staticconstexpr

Initial value:

=

Problem::kPadHeadDimQ

◆ kPadHeadDimV

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

bool ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::kPadHeadDimV

staticconstexpr

Initial value:

=

Problem::kPadHeadDimV

◆ kPadSeqLenK

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

bool ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::kPadSeqLenK = Problem::kPadSeqLenK

staticconstexpr

◆ kPadSeqLenQ

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

bool ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::kPadSeqLenQ = Problem::kPadSeqLenQ

staticconstexpr

◆ kQKHeaddim

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

index_t ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::kQKHeaddim = BlockFmhaShape::kQKHeaddim

staticconstexpr

◆ kQLoadOnce

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

bool ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::kQLoadOnce = true

staticconstexpr

◆ kStoreLSE

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

bool ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::kStoreLSE = Problem::kStoreLSE

staticconstexpr

◆ kSubQKHeaddim

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

index_t ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::kSubQKHeaddim = BlockFmhaShape::kSubQKHeaddim

staticconstexpr

◆ name

template<typename Problem_, typename Policy_ = BlockFmhaPipelineQRKSVSAsyncTrloadDefaultPolicy>

const char* ck_tile::BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ >::name = "qr_async_trload"

staticconstexpr

The documentation for this struct was generated from the following file:

block_fmha_pipeline_qr_ks_vs_async_trload.hpp

BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ > Struct Template Reference

BlockFmhaPipelineQRKSVSAsyncTrload&lt; Problem_, Policy_ &gt; Struct Template Reference#

Public Types

Public Member Functions

Static Public Member Functions

Static Public Attributes

Member Typedef Documentation

◆ AttentionVariant

◆ BiasDataType

◆ BlockFmhaShape

◆ FmhaMask

◆ KDataType

◆ LSEDataType

◆ OaccDataType

◆ ODataType

◆ PDataType

◆ Policy

◆ Problem

◆ QDataType

◆ RandValOutputDataType

◆ SaccDataType

◆ SMPLComputeDataType

◆ VDataType

◆ VLayout

Member Function Documentation

◆ GetSmemSize()

◆ operator()() [1/2]

◆ operator()() [2/2]

Member Data Documentation

◆ BiasEnum

◆ I0

◆ I1

◆ kAlignmentBias

◆ kAlignmentK

◆ kAlignmentOacc

◆ kAlignmentQ

◆ kAlignmentV

◆ kBlockPerCu

◆ kBlockSize

◆ kHasDropout

◆ kHasLogitsSoftCap

◆ kHasUnevenSplits

◆ kIsGroupMode

◆ kK0

◆ kK1

◆ kKLoadOnce

◆ kM0

◆ kN0

◆ kN1

◆ kNWarp

◆ kNXdl

◆ kPadHeadDimQ

◆ kPadHeadDimV

◆ kPadSeqLenK

◆ kPadSeqLenQ

◆ kQKHeaddim

◆ kQLoadOnce

◆ kStoreLSE

◆ kSubQKHeaddim

◆ name

BlockFmhaPipelineQRKSVSAsyncTrload< Problem_, Policy_ > Struct Template Reference#