GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer > Struct Template Reference

GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle&lt; ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer &gt; Struct Template Reference#

Composable Kernel: ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer > Struct Template Reference
Public Types | Static Public Member Functions | Static Public Attributes | List of all members
ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer > Struct Template Reference

#include <gridwise_gemm_multiple_d_welford_first_half_xdl_cshuffle.hpp>

Public Types

using ThisThreadBlock = ThisThreadBlock<BlockSize>
using GridwiseGemmPipe
using DefaultAGridDesc_AK0_M_AK1
using DefaultBGridDesc_BK0_N_BK1
using EGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock
using MeanVarGridDescriptor_MBlock_MPerBlock_NBlock
using CountGridDescriptor_MBlock_MPerBlock_NBlock
using DsGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock
using DefaultBlock2ETileMap
using DsGridPointer = decltype(MakeDsGridPointer())

Static Public Member Functions

__host__ static __device__ constexpr auto GetABlockDescriptor_AK0PerBlock_MPerBlock_AK1 ()
__host__ static __device__ constexpr auto GetBBlockDescriptor_BK0PerBlock_NPerBlock_BK1 ()
__host__ static __device__ constexpr auto GetCShuffleBlockDescriptor_MBlock_MPerBlock_NBlock_NPerBlock ()
static constexpr auto MakeDsGridPointer ()
__host__ static __device__ constexpr index_t GetSharedMemoryNumberOfByte ()
__host__ static __device__ constexpr auto MakeDefaultAGridDescriptor_AK0_M_AK1 (const AGridDesc_M_K &a_grid_desc_m_k)
__host__ static __device__ constexpr auto MakeDefaultBGridDescriptor_BK0_N_BK1 (const BGridDesc_N_K &b_grid_desc_n_k)
template<typename EGridDescriptor_M_N>
__host__ static __device__ constexpr auto MakeEGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock (const EGridDescriptor_M_N &e_grid_desc_m_n)
template<typename DsGridDescriptor_M_N>
__host__ static __device__ constexpr auto MakeDsGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock (const DsGridDescriptor_M_N &ds_grid_desc_m_n)
template<typename GridDescriptor_M_N>
__host__ static __device__ constexpr auto MakeMeanVarCountGridDescriptor_MBlock_MPerBlock_NBlock (const GridDescriptor_M_N &grid_desc_m_n)
__host__ static __device__ constexpr auto MakeDefaultBlock2ETileMap (const EGridDesc_M_N &e_grid_desc_m_n)
template<typename Block2ETileMap>
__host__ static __device__ constexpr bool CheckValidity (const AGridDesc_M_K &a_grid_desc_m_k, const BGridDesc_N_K &b_grid_desc_n_k, const DsGridDesc_M_N &ds_grid_desc_m_n, const EGridDesc_M_N &e_grid_desc_m_n, const Block2ETileMap &block_2_etile_map)
__host__ static __device__ constexpr bool CalculateHasMainKBlockLoop (index_t K)
template<bool HasMainKBlockLoop, typename AGridDesc_AK0_M_AK1, typename BGridDesc_BK0_N_BK1, typename Block2ETileMap>
static __device__ void Run (const ABDataType *__restrict__ p_a_grid, const ABDataType *__restrict__ p_b_grid, DsGridPointer p_ds_grid, EMeanVarDataType *__restrict__ p_e_grid, EMeanVarDataType *__restrict__ p_welford_mean_grid, EMeanVarDataType *__restrict__ p_welford_var_grid, int32_t *__restrict__ p_welford_count, void *__restrict__ p_shared, const AElementwiseOperation &a_element_op, const BElementwiseOperation &b_element_op, const CDEElementwiseOperation &cde_element_op, const AGridDesc_AK0_M_AK1 &a_grid_desc_ak0_m_ak1, const BGridDesc_BK0_N_BK1 &b_grid_desc_bk0_n_bk1, const DsGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock &ds_grid_desc_mblock_mperblock_nblock_nperblock, const EGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock &e_grid_desc_mblock_mperblock_nblock_nperblock, const MeanVarGridDescriptor_MBlock_MPerBlock_NBlock &mean_var_grid_desc_mblock_mperblock_nblock, const CountGridDescriptor_MBlock_MPerBlock_NBlock &count_grid_desc_mblock_mperblock_nblock, const Block2ETileMap &block_2_etile_map, index_t NRaw)

Static Public Attributes

static constexpr index_t NumDTensor = DsDataType::Size()
static constexpr auto I0 = Number<0>{}
static constexpr auto I1 = Number<1>{}
static constexpr auto I2 = Number<2>{}
static constexpr auto I3 = Number<3>{}
static constexpr auto I4 = Number<4>{}
static constexpr auto I5 = Number<5>{}
static constexpr auto I6 = Number<6>{}
static constexpr auto I7 = Number<7>{}
static constexpr auto AK1 = Number<AK1Value>{}
static constexpr auto BK1 = Number<BK1Value>{}
static constexpr auto AK0PerBlock = Number<KPerBlock / AK1Value>{}
static constexpr auto BK0PerBlock = Number<KPerBlock / BK1Value>{}

Member Typedef Documentation

◆ CountGridDescriptor_MBlock_MPerBlock_NBlock

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
using ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::CountGridDescriptor_MBlock_MPerBlock_NBlock
Initial value:
remove_cvref_t<decltype(MakeMeanVarCountGridDescriptor_MBlock_MPerBlock_NBlock(
CountGridDesc_M_NBlock{}))>
ck::remove_cvref_t
remove_cv_t< remove_reference_t< T > > remove_cvref_t
Definition type.hpp:297
ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle::MakeMeanVarCountGridDescriptor_MBlock_MPerBlock_NBlock
__host__ static __device__ constexpr auto MakeMeanVarCountGridDescriptor_MBlock_MPerBlock_NBlock(const GridDescriptor_M_N &grid_desc_m_n)
Definition gridwise_gemm_multiple_d_welford_first_half_xdl_cshuffle.hpp:246

◆ DefaultAGridDesc_AK0_M_AK1

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
using ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::DefaultAGridDesc_AK0_M_AK1
Initial value:
remove_cvref_t<decltype(MakeDefaultAGridDescriptor_AK0_M_AK1(AGridDesc_M_K{}))>
ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle::MakeDefaultAGridDescriptor_AK0_M_AK1
__host__ static __device__ constexpr auto MakeDefaultAGridDescriptor_AK0_M_AK1(const AGridDesc_M_K &a_grid_desc_m_k)
Definition gridwise_gemm_multiple_d_welford_first_half_xdl_cshuffle.hpp:180

◆ DefaultBGridDesc_BK0_N_BK1

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
using ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::DefaultBGridDesc_BK0_N_BK1
Initial value:
remove_cvref_t<decltype(MakeDefaultBGridDescriptor_BK0_N_BK1(BGridDesc_N_K{}))>
ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle::MakeDefaultBGridDescriptor_BK0_N_BK1
__host__ static __device__ constexpr auto MakeDefaultBGridDescriptor_BK0_N_BK1(const BGridDesc_N_K &b_grid_desc_n_k)
Definition gridwise_gemm_multiple_d_welford_first_half_xdl_cshuffle.hpp:196

◆ DefaultBlock2ETileMap

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
using ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::DefaultBlock2ETileMap
Initial value:
remove_cvref_t<decltype(MakeDefaultBlock2ETileMap(EGridDesc_M_N{}))>
ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle::MakeDefaultBlock2ETileMap
__host__ static __device__ constexpr auto MakeDefaultBlock2ETileMap(const EGridDesc_M_N &e_grid_desc_m_n)
Definition gridwise_gemm_multiple_d_welford_first_half_xdl_cshuffle.hpp:264

◆ DsGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
using ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::DsGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock
Initial value:
remove_cvref_t<decltype(MakeDsGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock(
DsGridDesc_M_N{}))>
ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle::MakeDsGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock
__host__ static __device__ constexpr auto MakeDsGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock(const DsGridDescriptor_M_N &ds_grid_desc_m_n)
Definition gridwise_gemm_multiple_d_welford_first_half_xdl_cshuffle.hpp:234

◆ DsGridPointer

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
using ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::DsGridPointer = decltype(MakeDsGridPointer())

◆ EGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
using ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::EGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock
Initial value:
remove_cvref_t<decltype(MakeEGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock(
EGridDesc_M_N{}))>
ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle::MakeEGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock
__host__ static __device__ constexpr auto MakeEGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock(const EGridDescriptor_M_N &e_grid_desc_m_n)
Definition gridwise_gemm_multiple_d_welford_first_half_xdl_cshuffle.hpp:212

◆ GridwiseGemmPipe

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
using ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::GridwiseGemmPipe
Initial value:
remove_cvref_t<
decltype(GridwiseGemmPipeline_Selector<PipelineVer, NumGemmKPrefetchStage, LoopSched>())>
ck::GridwiseGemmPipeline_Selector
constexpr auto GridwiseGemmPipeline_Selector()
Definition gridwise_gemm_pipeline_selector.hpp:31

◆ MeanVarGridDescriptor_MBlock_MPerBlock_NBlock

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
using ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::MeanVarGridDescriptor_MBlock_MPerBlock_NBlock
Initial value:
remove_cvref_t<decltype(MakeMeanVarCountGridDescriptor_MBlock_MPerBlock_NBlock(
MeanVarGridDesc_M_NBlock{}))>

◆ ThisThreadBlock

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
using ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::ThisThreadBlock = ThisThreadBlock<BlockSize>

Member Function Documentation

◆ CalculateHasMainKBlockLoop()

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
__host__ static __device__ constexpr bool ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::CalculateHasMainKBlockLoop ( index_t K)
inlinestaticconstexpr

◆ CheckValidity()

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
template<typename Block2ETileMap>
__host__ static __device__ constexpr bool ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::CheckValidity ( const AGridDesc_M_K & a_grid_desc_m_k,
const BGridDesc_N_K & b_grid_desc_n_k,
const DsGridDesc_M_N & ds_grid_desc_m_n,
const EGridDesc_M_N & e_grid_desc_m_n,
const Block2ETileMap & block_2_etile_map )
inlinestaticconstexpr

◆ GetABlockDescriptor_AK0PerBlock_MPerBlock_AK1()

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
__host__ static __device__ constexpr auto ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::GetABlockDescriptor_AK0PerBlock_MPerBlock_AK1 ( )
inlinestaticconstexpr

◆ GetBBlockDescriptor_BK0PerBlock_NPerBlock_BK1()

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
__host__ static __device__ constexpr auto ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::GetBBlockDescriptor_BK0PerBlock_NPerBlock_BK1 ( )
inlinestaticconstexpr

◆ GetCShuffleBlockDescriptor_MBlock_MPerBlock_NBlock_NPerBlock()

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
__host__ static __device__ constexpr auto ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::GetCShuffleBlockDescriptor_MBlock_MPerBlock_NBlock_NPerBlock ( )
inlinestaticconstexpr

◆ GetSharedMemoryNumberOfByte()

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
__host__ static __device__ constexpr index_t ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::GetSharedMemoryNumberOfByte ( )
inlinestaticconstexpr

◆ MakeDefaultAGridDescriptor_AK0_M_AK1()

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
__host__ static __device__ constexpr auto ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::MakeDefaultAGridDescriptor_AK0_M_AK1 ( const AGridDesc_M_K & a_grid_desc_m_k)
inlinestaticconstexpr

◆ MakeDefaultBGridDescriptor_BK0_N_BK1()

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
__host__ static __device__ constexpr auto ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::MakeDefaultBGridDescriptor_BK0_N_BK1 ( const BGridDesc_N_K & b_grid_desc_n_k)
inlinestaticconstexpr

◆ MakeDefaultBlock2ETileMap()

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
__host__ static __device__ constexpr auto ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::MakeDefaultBlock2ETileMap ( const EGridDesc_M_N & e_grid_desc_m_n)
inlinestaticconstexpr

◆ MakeDsGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock()

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
template<typename DsGridDescriptor_M_N>
__host__ static __device__ constexpr auto ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::MakeDsGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock ( const DsGridDescriptor_M_N & ds_grid_desc_m_n)
inlinestaticconstexpr

◆ MakeDsGridPointer()

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
constexpr auto ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::MakeDsGridPointer ( )
inlinestaticconstexpr

◆ MakeEGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock()

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
template<typename EGridDescriptor_M_N>
__host__ static __device__ constexpr auto ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::MakeEGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock ( const EGridDescriptor_M_N & e_grid_desc_m_n)
inlinestaticconstexpr

◆ MakeMeanVarCountGridDescriptor_MBlock_MPerBlock_NBlock()

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
template<typename GridDescriptor_M_N>
__host__ static __device__ constexpr auto ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::MakeMeanVarCountGridDescriptor_MBlock_MPerBlock_NBlock ( const GridDescriptor_M_N & grid_desc_m_n)
inlinestaticconstexpr

◆ Run()

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
template<bool HasMainKBlockLoop, typename AGridDesc_AK0_M_AK1, typename BGridDesc_BK0_N_BK1, typename Block2ETileMap>
__device__ void ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::Run ( const ABDataType *__restrict__ p_a_grid,
const ABDataType *__restrict__ p_b_grid,
DsGridPointer p_ds_grid,
EMeanVarDataType *__restrict__ p_e_grid,
EMeanVarDataType *__restrict__ p_welford_mean_grid,
EMeanVarDataType *__restrict__ p_welford_var_grid,
int32_t *__restrict__ p_welford_count,
void *__restrict__ p_shared,
const AElementwiseOperation & a_element_op,
const BElementwiseOperation & b_element_op,
const CDEElementwiseOperation & cde_element_op,
const AGridDesc_AK0_M_AK1 & a_grid_desc_ak0_m_ak1,
const BGridDesc_BK0_N_BK1 & b_grid_desc_bk0_n_bk1,
const DsGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock & ds_grid_desc_mblock_mperblock_nblock_nperblock,
const EGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock & e_grid_desc_mblock_mperblock_nblock_nperblock,
const MeanVarGridDescriptor_MBlock_MPerBlock_NBlock & mean_var_grid_desc_mblock_mperblock_nblock,
const CountGridDescriptor_MBlock_MPerBlock_NBlock & count_grid_desc_mblock_mperblock_nblock,
const Block2ETileMap & block_2_etile_map,
index_t NRaw )
inlinestatic

Member Data Documentation

◆ AK0PerBlock

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
auto ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::AK0PerBlock = Number<KPerBlock / AK1Value>{}
staticconstexpr

◆ AK1

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
auto ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::AK1 = Number<AK1Value>{}
staticconstexpr

◆ BK0PerBlock

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
auto ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::BK0PerBlock = Number<KPerBlock / BK1Value>{}
staticconstexpr

◆ BK1

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
auto ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::BK1 = Number<BK1Value>{}
staticconstexpr

◆ I0

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
auto ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::I0 = Number<0>{}
staticconstexpr

◆ I1

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
auto ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::I1 = Number<1>{}
staticconstexpr

◆ I2

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
auto ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::I2 = Number<2>{}
staticconstexpr

◆ I3

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
auto ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::I3 = Number<3>{}
staticconstexpr

◆ I4

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
auto ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::I4 = Number<4>{}
staticconstexpr

◆ I5

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
auto ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::I5 = Number<5>{}
staticconstexpr

◆ I6

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
auto ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::I6 = Number<6>{}
staticconstexpr

◆ I7

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
auto ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::I7 = Number<7>{}
staticconstexpr

◆ NumDTensor

template<typename ABDataType, typename AccDataType, typename CShuffleDataType, typename DsDataType, typename EMeanVarDataType, typename AElementwiseOperation, typename BElementwiseOperation, typename CDEElementwiseOperation, InMemoryDataOperationEnum EGlobalMemoryDataOperation, typename AGridDesc_M_K, typename BGridDesc_N_K, typename DsGridDesc_M_N, typename EGridDesc_M_N, typename MeanVarGridDesc_M_NBlock, typename CountGridDesc_M_NBlock, index_t NumGemmKPrefetchStage, index_t BlockSize, index_t MPerBlock, index_t NPerBlock, index_t KPerBlock, index_t AK1Value, index_t BK1Value, index_t MPerXdl, index_t NPerXdl, index_t MXdlPerWave, index_t NXdlPerWave, typename ABlockTransferThreadClusterLengths_AK0_M_AK1, typename ABlockTransferThreadClusterArrangeOrder, typename ABlockTransferSrcAccessOrder, index_t ABlockTransferSrcVectorDim, index_t ABlockTransferSrcScalarPerVector, index_t ABlockTransferDstScalarPerVector_AK1, bool AThreadTransferSrcResetCoordinateAfterRun, index_t ABlockLdsExtraM, typename BBlockTransferThreadClusterLengths_BK0_N_BK1, typename BBlockTransferThreadClusterArrangeOrder, typename BBlockTransferSrcAccessOrder, index_t BBlockTransferSrcVectorDim, index_t BBlockTransferSrcScalarPerVector, index_t BBlockTransferDstScalarPerVector_BK1, bool BThreadTransferSrcResetCoordinateAfterRun, index_t BBlockLdsExtraN, index_t CShuffleMXdlPerWavePerShuffle, index_t CShuffleNXdlPerWavePerShuffle, typename PostShuffleThreadClusterSize_M_N, index_t PostShuffleScalarPerVector, LoopScheduler LoopSched, PipelineVersion PipelineVer = PipelineVersion::v1>
index_t ck::GridwiseGemmMultipleDWelfordFirstHalf_xdl_cshuffle< ABDataType, AccDataType, CShuffleDataType, DsDataType, EMeanVarDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, EGlobalMemoryDataOperation, AGridDesc_M_K, BGridDesc_N_K, DsGridDesc_M_N, EGridDesc_M_N, MeanVarGridDesc_M_NBlock, CountGridDesc_M_NBlock, NumGemmKPrefetchStage, BlockSize, MPerBlock, NPerBlock, KPerBlock, AK1Value, BK1Value, MPerXdl, NPerXdl, MXdlPerWave, NXdlPerWave, ABlockTransferThreadClusterLengths_AK0_M_AK1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_AK1, AThreadTransferSrcResetCoordinateAfterRun, ABlockLdsExtraM, BBlockTransferThreadClusterLengths_BK0_N_BK1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_BK1, BThreadTransferSrcResetCoordinateAfterRun, BBlockLdsExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, PostShuffleThreadClusterSize_M_N, PostShuffleScalarPerVector, LoopSched, PipelineVer >::NumDTensor = DsDataType::Size()
staticconstexpr
The documentation for this struct was generated from the following file: