/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/fused_moe/kernel/moe_sorting_kernel.hpp File Reference

• develop
• include
• ck_tile
• ops
• fused_moe
• kernel

Classes | Namespaces | Macros | Functions

moe_sorting_kernel.hpp File Reference

#include "ck_tile/core.hpp"
#include "ck_tile/ops/common.hpp"
#include "ck_tile/ops/elementwise.hpp"
#include "ck_tile/host/hip_check_error.hpp"
#include <string>
#include <type_traits>

Go to the source code of this file.

Classes
struct	ck_tile::MoeSortingHostArgs
struct	ck_tile::MoeSortingKernel< Problem_ >
struct	ck_tile::MoeSortingKernel< Problem_ >::Kargs
struct	ck_tile::MoeSortingKernel< Problem_ >::simple_smem_indexer
struct	ck_tile::MoeSortingClearWorkspaceKernel< Problem_ >
struct	ck_tile::MoeSortingClearWorkspaceKernel< Problem_ >::Kargs
struct	ck_tile::MoeSortingMultiPhaseKernel_P0< Problem_ >
struct	ck_tile::MoeSortingMultiPhaseKernel_P0< Problem_ >::Kargs
struct	ck_tile::MoeSortingMultiPhaseKernel_P1< Problem_ >
struct	ck_tile::MoeSortingMultiPhaseKernel_P1< Problem_ >::Kargs
struct	ck_tile::MoeSortingMultiPhaseKernel_P2< Problem_ >
struct	ck_tile::MoeSortingMultiPhaseKernel_P2< Problem_ >::Kargs
struct	ck_tile::MoeSortingMultiPhaseKernel_P3< Problem_ >
struct	ck_tile::MoeSortingMultiPhaseKernel_P3< Problem_ >::Kargs
struct	ck_tile::MoeSortingMultiPhaseKernel_P23< Problem_ >
struct	ck_tile::MoeSortingMultiPhaseKernel_P23< Problem_ >::Kargs

Namespaces
	ck_tile
	ck_tile::impl

Macros
#define	MOE_SORTING_MOCK_ID(token_id_, topk_id_)    static_cast<uint32_t>(((token_id_) & 0x00ffffff) | (((topk_id_) & 0xff) << 24))
#define	MOE_SORTING_USE_EX_KERNEL   1
#define	MOE_SORTING_FUSE_MP_01   0
#define	MOE_SORTING_FMOE_2D_BUF   1

Functions
constexpr CK_TILE_HOST auto	ck_tile::moe_sorting_get_smem_row_col (int tokens_, int num_experts_)
CK_TILE_HOST index_t	ck_tile::moe_sorting_get_sub_token (int tokens_, int num_experts_)
CK_TILE_HOST_DEVICE index_t	ck_tile::impl::moe_sorting_mp_mesh_stride (index_t tokens)
CK_TILE_HOST index_t	ck_tile::impl::moe_sorting_mesh_byte_size (index_t tokens_, index_t, index_t topk_)
CK_TILE_HOST_DEVICE index_t	ck_tile::impl::moe_sorting_mp_mesh_smem_size (index_t tokens, index_t num_experts, index_t topk)
CK_TILE_HOST_DEVICE index_t	ck_tile::impl::moe_sorting_mp_cumsum_smem_size (index_t num_experts)
CK_TILE_HOST_DEVICE index_t	ck_tile::impl::moe_sorting_mp_sem_smem_size ()
template<typename T , typename F , index_t wave_size_ = get_warp_size()>
constexpr CK_TILE_DEVICE T	ck_tile::impl::moe_sorting_wave_reduce (T local, F reduce_f, number< wave_size_ >={})
template<typename data_t , int wave_size>
CK_TILE_DEVICE void	ck_tile::impl::moe_sorting_wave_cumsum (data_t &thread_data)
template<index_t BLOCK_SIZE = 256>
CK_TILE_DEVICE void	ck_tile::impl::moe_buf_set_zero_kernel (uint8x16_t *buf, long_index_t buf_bytes, index_t gid)
template<index_t BLOCK_SIZE = 256>
CK_TILE_DEVICE void	ck_tile::impl::moe_buf_set_zero_kernel_2d (void *buf, index_t row, index_t col, index_t elem_bytes, index_t gid, index_t blocks)
CK_TILE_HOST bool	ck_tile::moe_sorting_is_oneshot (int tokens_, int num_experts_)
CK_TILE_HOST index_t	ck_tile::moe_sorting_mp_get_workspace_size (int tokens_, int num_experts_, int topk_)
CK_TILE_HOST index_t	ck_tile::moe_sorting_get_workspace_size (int tokens_, int num_experts_, int topk_, int dispatch_policy_)
constexpr CK_TILE_HOST auto	ck_tile::impl::moe_sorting_get_smem_size_p23 (int num_experts_)

Macro Definition Documentation

◆ MOE_SORTING_FMOE_2D_BUF

#define MOE_SORTING_FMOE_2D_BUF   1

◆ MOE_SORTING_FUSE_MP_01

#define MOE_SORTING_FUSE_MP_01   0

◆ MOE_SORTING_MOCK_ID

#define MOE_SORTING_MOCK_ID	(		token_id_,
			topk_id_
	)		static_cast<uint32_t>(((token_id_) & 0x00ffffff) | (((topk_id_) & 0xff) << 24))

◆ MOE_SORTING_USE_EX_KERNEL

#define MOE_SORTING_USE_EX_KERNEL   1