MegaMoE 在 MoE 系统谱系中的位置：从 DeepSpeed-MoE 到 CCFuser 的七方对比

MoE 的稀疏 FLOPs 不会自动变成推理加速。DeepSpeed-MoE 最早把这件事讲清楚：小 batch 和 decode 阶段，瓶颈是权重带宽和通信，不是算力。此后的系统工作基本都在回答同一个问题的不同侧面——动态路由带来的不规则性，该在哪一层吸收？

本文以 DeepSeek DeepGEMM 里的 fp8_fp4_mega_moe 单 kernel 实现为坐标，把六项相关工作放在一起对照。代码行号对应 2026-09-08 的 DeepGEMM @559d79f。

---

一、谱系：每一步继承了什么问题

GShard / Switch Transformer (2020-2021)
  │ 静态 capacity factor 的两难：丢 token（掉点）或补零（浪费算力）
  ▼
DeepSpeed-MoE (2022)
  ├─ 解决：推理时参数显存与访存瓶颈
  ├─ 论点：MoE 推理是访存带宽受限，不是 FLOP 受限
  └─ 遗留：仍依赖静态 capacity factor 与稠密 Einsum 布局
      ▼
FasterMoE (PPoPP 2022)
  ├─ 解决：动态路由倾斜造成的 straggler、集合通信阻塞
  ├─ 手段：dynamic shadowing 复制热门专家；smart scheduling 跨 stream 流水
  └─ 遗留：影子权重吃显存；仍在传统 batched GEMM + capacity factor 框架内
      ▼
Tutel (MLSys 2023)
  ├─ 解决：静态并行策略不灵活；千卡规模下 flat all-to-all 扩展性
  ├─ 手段：运行时零成本 DP/EP/MP 切换；4 阶段 2DH 集合通信；动态容量不丢 token
  └─ 遗留：容量取 max(T_e) 导致所有欠载专家补零，浪费 2–11 倍算力
      ▼
MegaBlocks (MLSys 2023)
  ├─ 解决：补零与丢 token 的两难
  ├─ 手段：MoE 表达成块稀疏矩阵乘（BCSR-BCOO，128×128 块，SDD/DSD）
  │        padding 浪费上界收紧到每专家 ≤ 127 token
  └─ 遗留：仍需在 HBM 里 materialize 置换后的输入张量和块稀疏中间激活
      ▼
ScatterMoE (2024)
  ├─ 解决：MegaBlocks 的 HBM 分配与冗余拷贝
  ├─ 手段：ParallelLinear（Triton）把 grouped GEMM 与间接寻址融合，
  │        HBM↔SRAM 之间直接 gather/scatter，不落中间缓冲
  └─ 遗留：单机；无多机通信原语；无单 token 解码专用 kernel
      ▼
      ├──────────────────────────────┐
      ▼                              ▼
CCFuser (PPoPP 2025)          DeepGEMM Mega MoE (2026)
分阶段融合 + NVSHMEM            单 persistent kernel + symmetric memory

最后一步值得强调：ScatterMoE 自己承认的两条局限，正是 Mega MoE 的起点。

ScatterMoE is currently single-node focused and does not provide multi-node communication primitives… It also lacks a specialized decoding kernel for single-token autoregressive generation.

Mega MoE 对这两条的回应分别是：把跨卡 dispatch/combine 直接写进 kernel；以及把 tile 大小做成 serving regime 的函数，最小档明确标注给 decode 和 RL rollout 用。

---

二、Mega MoE 的具体实现

2.1 一个 kernel 里的六种 warp 角色

deep_gemm/include/deep_gemm/impls/sm100_fp8_fp4_mega_moe.cuh:329 起，CTA 内的 warp 按 index 分成互不相同的角色：

warp 范围	角色	行号
[0, kNumDispatchWarps)	跨卡 dispatch / combine 搬运	:329
kNumDispatchWarps + 0	激活 + SFA 的 TMA load	:669
kNumDispatchWarps + 1	权重 + SFB 的 TMA load	:735
kNumDispatchWarps + 2	MMA 发射（tcgen05，仅 leader CTA）	:794
kNumDispatchWarps + 3	调度器 mainloop（仅 leader CTA）	:920
>= kNumDispatchWarps + kNumMMANonEpilogueWarps	epilogue：TMEM → SwiGLU + amax + 量化 → 存回 / combine 写回	:928

每个角色进入分支后先调整寄存器配额（warpgroup_reg_dealloc / warpgroup_reg_alloc），静态断言总量 ≤ 64512（:319）。

对比意义：FasterMoE 的 fine-grained scheduling 和 Tutel 的 adaptive pipelining 都是在 kernel 之上编排多个 kernel 与 stream；Mega MoE 把重叠下沉到单 kernel 内的 warp 特化，靠硬件 warp 调度而非 CUDA stream。

2.2 pull 模型 + rank 轮转交错

dispatch warp 遍历的不是本地 token，而是本地专家池的目标槽位，反推每个槽位该从哪个 rank 拉哪个 token：

// Pull token data and SF from remote ranks into local L1 buffer   (:414)
// Round-robin rank selection via iterative min-peeling            (:461)

两点含义：

1. 没有集合通信。不做 all-to-all，消费端直接去对端符号内存取数。DeepSpeed-MoE 的分层 all-to-all、Tutel 的 2DH 优化的都是集合通信本身，这里整个绕开了。
2. rank 轮转交错。不是 rank0 的 token 排完再排 rank1，而是让每个 BLOCK_M tile 内的 token 尽量来自不同 rank，把 NVLink 读流量摊平。这是对负载不均的回应，作用在数据布局层，而非 FasterMoE 的专家复制层。

2.3 中间激活是有界环形缓冲

这是与 MegaBlocks / ScatterMoE / CCFuser 差别最大的一点。csrc/apis/mega.hpp:57-65：

int num_ring_tokens = 0;
for (const auto& block_m: layout::kCandidateBlockM) {
    const auto num_pool_blocks = ceil_div(num_max_routed_tokens, block_m) + num_experts_per_rank;
    const auto num_live_pool_blocks = sched::get_num_max_live_pool_blocks(
        num_pool_blocks, num_sms, hidden, intermediate_hidden);
    num_ring_tokens = std::max(num_ring_tokens, num_live_pool_blocks * block_m);
}
num_ring_tokens = math::align(num_ring_tokens, layout::kLCMCandidateBlockM);

l1_acts 的第一维是 num_ring_tokens——按「同时在飞的 block 数」定尺寸，不是按路由后的总 token 数。

准确说明

这个环形缓冲位于符号内存（全局显存），不是 SRAM。它的意义不是「不落显存」，而是容量与路由后的总 token 数解耦：ScatterMoE 消除的是置换缓冲，Mega MoE 进一步让 L1 输出也不必按总量分配。

dispatch → L1 → L2 → combine 四段用四组计数器在环上做生产者–消费者同步：

get_l1_full_count_ptr / get_l1_empty_count_ptr
get_l2_full_count_ptr / get_l2_empty_count_ptr

消费端自旋等待（:527-530、:990），环块索引一律取模 kNumRingBlocks。

2.4 routing index 直接驱动跨卡写地址

struct TokenSrcMetadata {
    uint32_t rank_idx;
    uint32_t token_idx;
    uint32_t topk_idx;
};

dispatch 时为每个池内 token 写下来源（:587），combine 时 epilogue warp 读它算出对端地址，经 NVLink 直接写入目标 rank 的缓冲（:1280、:1294）。

这正是 ScatterMoE「routing index 驱动 GEMM 寻址」的思路，但 scatter 的目标是另一块 GPU 的显存。ScatterMoE 的 gather/scatter 在单卡 HBM↔SRAM 之间；Mega MoE 把同一个机制拉到了 NVLink 域。

2.5 块大小按 serving regime 选档

候选集（deep_gemm/include/deep_gemm/layout/mega_moe.cuh:11）：

static constexpr int kCandidateBlockM[7] = {8, 16, 32, 64, 96, 128, 192};

选择逻辑在 csrc/jit_kernels/heuristics/mega_moe.hpp:80 起，判据是 $\text{num\_expected\_tokens\_per\_expert} = \frac{\text{num\_tokens} \times \text{num\_ranks} \times \text{num\_topk}}{\text{num\_experts}}$

每专家期望 token	BLOCK_M	BLOCK_K	代码原注释
≤ 8.5	16	256	Really small token-per-expert (e.g. RL long-tail rollout)
≤ 16.5	32	128	Small batch size, small EP, decoding, e.g. 6/384 experts, EP8, bsz 128
≤ 32.5	64	128	Medium batch size, small EP, decoding, e.g. 6/384 experts, EP8, bsz 256
≤ 64.5	96	128	Large batch size, small EP, decoding, e.g. 6/384 experts, EP8, bsz 512
≤ 96.5	128	128	Medium batch size, Medium EP, decoding, EP16 bsz 256, or EP32 bsz 128
其余	192	128	Prefill, or large EP decoding

MegaBlocks 用固定 128×128 块加块稀疏消除 padding，并承认小 batch 或 $T_e < 128$ 时 fringe 严重；Mega MoE 保留稠密 tile，但把块大小做成 serving regime 的函数，最小档到 16。注释直接点名 RL rollout 与 decode——它是从推理场景倒推的设计。

2.6 dropless，但方式不同

池容量按最坏情况开（layout/mega_moe.cuh:18）：

$\text{pool} = \text{num\_ranks} \times \text{num\_max\_tokens\_per\_rank} \times \min(\text{topk}, \text{experts\_per\_rank}) + \text{experts\_per\_rank} \times 191$

对齐到 384。没有 capacity factor、不丢 token——代价是显存按最坏情况预留，而不是像 MegaBlocks 那样用稀疏表示避开。

2.7 共享专家是同一队列里的两个 phase

enum class BlockPhase { None, Linear1, Linear2, SharedLinear1, SharedLinear2 };

调度器是全局原子计数器驱动的工作队列（get_next_task_idx 用 atomic_add），L2 任务会自旋等待对应的 L1 任务先被取走：

const auto num_required_l1_tasks = (task_info.pool_block_idx + 1) * kNumL1Clusters;
while (ptx::ld_volatile(workspace.get_l1_task_count_ptr()) < num_required_l1_tasks) {}

依赖关系用 kernel 内的计数器表达，而不是靠 kernel 边界。 共享专家的 GEMM 与路由专家的 GEMM 在同一个任务流里被调度。

---

三、CCFuser：最接近的对照组

Hulin Wang, Yaqi Xia, Donglin Yang, Xiaobo Zhou, Dazhao Cheng. Harnessing Inter-GPU Shared Memory for Seamless MoE Communication-Computation Fusion. PPoPP ‘25. DOI 10.1145/3710848.3710868 · 代码 WongHuLin/CCFuser（无 arXiv 预印本）

CCFuser 和 Mega MoE 的目标几乎相同——把 MoE 通信融进计算——但四个关键选择都不一样。

相同点

• 都用「跨 GPU 共享内存」而非集合通信：CCFuser 走 NVSHMEM PGAS 对称堆（nvshmem_malloc），Mega MoE 走 PyTorch symmetric memory 的对等映射指针
• dispatch 都是 pull：CCFuser 用 nvshmemx_get16_nbi_block，Mega MoE 由消费端反推源 rank
• 都在设备侧发起通信，不经主机

不同点

维度	CCFuser	Mega MoE
kernel 结构	分阶段融合，多次 launch：local scatter → barrier → 融合 dispatch+GEMM1 → 激活（PyTorch）→ 融合 GEMM2+combine → barrier → gather	单 persistent kernel，整层一次 launch
内存抽象	NVSHMEM PGAS，可走 UCX/IB 出节点	symmetric memory 对等指针，纯 NVLink，无 RDMA，kNumMaxRanks = 72
通信方向	混合：dispatch pull，combine push（nvshmemx_put128_nbi_block）	全 pull
中间激活	完整落 HBM：merged_input、ffn1_out 用 torch::zeros 分配，供 autograd 保存	有界环形缓冲，容量与总 token 数解耦
CTA 内模型	同构协作 warp，全部 warp 一起发 async 传输再一起算 WMMA	warp 特化，六种角色
融合粒度	token block，固定 $BM=64$	MMA tile / warp tile，$BM \in \{16,32,64,96,128,192\}$ 动态选
硬件	A100，Ampere WMMA，FP16	SM100，tcgen05 + TMA，FP8×FP4
目标负载	分布式训练（含反向与权重梯度 kernel）	推理（Mega MoE 无反向）
对比基线 / 加速比	FastMoE 2.96×、FasterMoE 2.48×，峰值 4.34×	见 DeepGEMM PR #316

CCFuser 自己列的局限里，有两条正好被 Mega MoE 的设计绕开：

1. 固定 $BM=64$ 离散化：token 数不是 64 的倍数就要补齐，负载倾斜时碎片明显 → Mega MoE 用六档动态块大小
2. 本地/远程块配对依赖：靠「一个本地块配一个远程块」在 CTA 内互相掩藏延迟，若某 GPU 本地块过少就无法掩藏，需要离线的 expert reassignment 图变换 → Mega MoE 的 warp 特化让 dispatch warp 与 MMA warp 天然独立，不依赖配对

另外两条是它承认的代价：中间激活全落 HBM 吃带宽；多次 kernel launch 与 stream barrier 把 CPU 留在关键路径上。

---

四、七方对比矩阵

维度	DeepSpeed-MoE	FasterMoE	Tutel	MegaBlocks	ScatterMoE	CCFuser	Mega MoE
年份/会议	2022	PPoPP’22	MLSys’23	MLSys’23	2024	PPoPP’25	2026
主要目标	推理服务扩展	训练负载不均	统一自适应并行	dropless 训练	消除 HBM 拷贝	通信计算融合	推理单核融合
丢 token	是（静态容量）	是	可配（补零式 dropless）	否	否	否	否（最坏情况预留）
GEMM 原语	稠密 batched	稠密 batched	稠密 batched	块稀疏 SDD/DSD	融合 grouped GEMM	WMMA 分块	tcgen05 分块
应对倾斜	PR-MoE 架构	复制热门专家	调 $r$ + 补零	每专家动态块数	指针偏移动态 tile	离线专家重分配	rank 轮转交错
通信	分层 all-to-all	点对点流水	2DH 集合	标准 EP	无（单机）	NVSHMEM get/put	symmetric memory pull
中间激活	落盘	落盘	落盘（补零膨胀）	落盘	不 materialize 置换缓冲	落盘	有界环
融合层次	库调用	跨 stream 编排	跨 stream 编排	kernel 内块稀疏	kernel 内 gather/scatter	分阶段跨 kernel	单 kernel 内 warp 特化
规模域	scale-out	scale-out	scale-out	scale-out	单机	两者皆可	仅 scale-up（NVLink）
训练/推理	两者	训练	两者	训练	两者	训练	仅推理前向

---

五、Mega MoE 放弃了什么

对比之后，代价也清楚了：

• 没有 scale-out。全仓库 grep 不到任何 RDMA / IBGDA / NVSHMEM，kNumMaxRanks = 72 对着 NVL72。跨节点仍得回到 DeepEP。CCFuser 和 Tutel 在这一维上更通用。
• 没有反向。csrc/apis/mega.hpp 只导出 fp8_fp4_mega_moe 和 bf16_mega_moe 两个前向入口，impls/ 下没有反向大 kernel。训练路径仍是「DeepEP dispatch → DeepGEMM grouped GEMM → TileKernels SwiGLU → GEMM → combine」的分立流水。
• 放弃了自适应并行。Tutel 用参数 $r$ 在激活通信与参数通信之间搜最优点；Mega MoE 把「搬 token」这个选择写死了，不复制专家、不切换 DP/EP/MP。
• 放弃了稀疏表示。MegaBlocks 用 BCSR-BCOO 从根上消除 padding；Mega MoE 仍是稠密 tile，只是把块大小选细。

换来的是：整层一次 kernel launch，CPU 完全离开关键路径，通信与计算的重叠由硬件 warp 调度器完成，中间激活容量与路由总量解耦。这是一套彻底面向低延迟推理的取舍。

---

六、关于 UniMoE 的说明

交叉检索 DBLP、ACM DL、IEEE Xplore、arXiv 后确认：不存在名为 UniMoE / Uni-MoE 的系统或 GPU kernel 论文。

该名称下的文献是哈尔滨工业大学（深圳）与腾讯的多模态模型系列：

• Uni-MoE: Scaling Unified Multimodal LLMs with Mixture of Experts（arXiv:2405.11273，2024-05）
• Uni-MoE-2.0-Omni（arXiv:2511.12609）
• UniMoE-Audio（arXiv:2510.13344）

官方代码库 HITsz-TMG/Uni-MoE 中没有任何 .cu / C++ / Triton kernel，MoE 层是标准 PyTorch 线性层，跨卡通信直接用 DeepSpeed / PyTorch 的 all_to_all_single。它与本文讨论的 kernel 级工作不在同一层次。

---

七、存疑

未验证

1. Mega MoE 的性能数据未复现

本文所有关于 Mega MoE 的结论均来自源码阅读，未在 SM100 硬件上实测。官方性能数据见 DeepGEMM PR #316。

2. CCFuser 的多机表现

其仓库含 2 节点 8 卡的 MPI 启动脚本（NVSHMEM_REMOTE_TRANSPORT='ucx'），但论文主要结果集中在单节点 4×A100。设备端 nvshmemx_get16_nbi_block 跨 InfiniBand 的实际退化程度未能从公开日志确认。

3. FasterMoE 的部分细节

ACM DL 的 camera-ready PDF 抓取被拦截，roofline 模型的精确表述来自开源实现（thu-pacman/FasterMoE）与会议元数据推断。

---

版本

DeepGEMM 行号对应 2026-09-08 的本地检出 @559d79f。相关阅读：TileKernels 与 DeepGEMM 的分界线 · MoE 系统七篇精读 · MoE 的两个上游创新。