MoE 系统七篇精读：机制、代价与设计决策

MoE 的稀疏 FLOPs 不会自动变成加速。这条线上的七篇工作，每一篇都在回答前一篇留下的问题。本文逐篇拆机制，最后收成设计决策。

起点：为什么稀疏 FLOPs 不等于加速

DeepSpeed-MoE 把这件事量化了。自回归解码时 batch 小、step 长度为 1，算术强度极低，执行严格受显存带宽约束——延迟等于把参数从 HBM 读进寄存器的时间。

关键在于单 token 激活的参数量不等于一个 batch 激活的参数量：

1.3B+MoE-128 模型里，单个 token 只激活 1.3B 参数，但一个 batch 里的不同 token 合起来会激活到全部 52B。不做专家分布式部署去服务这个模型，每步要读进的参数字节数比质量相当的 6.7B 稠密模型多约 10×。

所以 FLOPs 省下来毫无用处——执行单元在等 DRAM。加速的唯一来源是靠 EP 把几十块 GPU 的显存带宽聚合起来，让每块卡只流式读自己那份专家切片。

这个论断是后面所有工作的前提。

一、DeepSpeed-MoE (ICML 2022)

Rajbhandari et al., arXiv:2201.05596

1.1 PR-MoE：两个经验现象

Pyramid（金字塔）：深层比浅层更受益于稀疏容量，浅层学的是通用的表示级特征。于是浅层放少的专家、深层放多的——例如前 10 层各 32 个专家，最后 2 层各 64 个。

Residual（残差）：Top-2 门控翻倍容量的同时也翻倍了 all-to-all 通信量。Residual-MoE 改成每个 token 过一个固定的稠密 MLP 加一个 Top-1 选出的稀疏专家，两者相加：

$\mathbf{y} = \text{MLP}_{\text{fixed}}(\mathbf{x}) + g_e(\mathbf{x}) \cdot \text{Expert}_e(\mathbf{x})$

稀疏专家扮演稠密基底之上的误差修正残差，用 Top-1 的通信量拿到 Top-2 的精度。两者合起来参数量降 3×。

1.2 MoS：分阶段蒸馏

把深的 PR-MoE 教师蒸馏成浅的学生（24 层 → 21 层）。作者发现全程蒸馏反而在训练后期掉点：学生容量已经缩了，还要同时优化匹配教师的 KD loss 和硬的语言建模交叉熵，导致欠拟合。

分阶段 KD 的做法是在 400K 步就停掉蒸馏损失，剩下的步数只用交叉熵训。保住 99.1%–99.5% 精度，额外再降 1.25× 参数（总计相对标准 MoE 降 3.7×）。

1.3 两个通信优化

分层 all-to-all：把跨 $p$ 卡的扁平 all-to-all 拆成四步——本地数据变换、节点内 NVLink all-to-all、布局转置、节点间 all-to-all。通信跳数从 $O(p)$ 降到 $O(G + p/G)$ 。

注意这里的取舍：总数据量翻倍，但小 batch 时延迟由每条消息的连接建立开销主导而非链路带宽，所以总延迟反而显著下降。

并行度协同通信：从 TP 的注意力层进到 EP 的专家层时，前面的 TP all-reduce 已经把数据在 $L$ 个张量并行 rank 之间复制了一遍。既然如此，all-to-all 就不必跨全部 $p$ 张卡，只在具有相同 TP rank 编号的 $p/L$ 个设备之间做即可。延迟从 $O(p)$ 降到 $O(p/L)$ ——128 卡 8 路 TP 下是 $(128C_1 + C_2) \to (16C_1 + C_2)$ 。

反方向（EP → TP）则是先在 $p/L$ 子集内 all-to-all，再在 $L$ 个 rank 之间做节点内 all-gather，界在 $O(p/L) + O(L)$ 。

1.4 结果与遗留

256 张 A100 上以 25ms 延迟服务 1.5T 参数 MoE；相比 PyTorch MoE 基线延迟降 7.3×，相比质量相当的稠密模型快 4.5×、便宜 9×。

遗留问题：仍依赖静态 capacity factor 和丢 token；算子层面还是稠密 einsum 布局。

二、FasterMoE (PPoPP 2022)

He et al., DOI:10.1145/3503221.3508418

2.1 Dynamic shadowing：搬 token 还是搬权重

标准 EP 下，所有路由到专家 $e$ 的 token 都得发到 $e$ 所在的 worker。路由一倾斜，那个 worker 就成了 straggler。

Dynamic shadowing 反过来搬权重：把热门专家的参数复制到别的 worker，让它们本地处理，不必把 token 发出去。

判据是每步在线算的成本模型。不做 shadowing 的延迟：

$T_{\text{baseline}} = \frac{12 B_w \alpha d_{\text{model}}^2 \cdot \text{sizeof}}{\text{BW}_{\text{mm}}} + \frac{4 B_w d_{\text{model}} \cdot \text{sizeof}}{\text{BW}_{\text{net}}}$

其中 $B_w = \max_j(\text{worker } j \text{ 上的 token 数})$ ， $\alpha = d_{\text{hidden}}/d_{\text{model}}$ 。

对最热的 $i$ 个专家做 shadowing 的延迟：

$T_{\text{shadow}}(i) = B_k(i) \cdot \tau_{\text{comp}} + i \cdot \frac{4 \alpha d_{\text{model}}^2 E \cdot \text{sizeof}}{\text{BW}_{\text{net}}}$

$B_k(i)$ 是剩下未被复制的专家里新的瓶颈负载，第二项是 $i$ 份参数广播同步的开销。

取 $i^* = \arg\min_i T_{\text{shadow}}(i)$ ，若 $T_{\text{shadow}}(i^*) < T_{\text{baseline}}$ 就对最热的 $i^*$ 个专家做 shadowing。

这个模型的形状值得记：第一项随 $i$ 递减（瓶颈变轻），第二项随 $i$ 递增（广播变多），所以存在最优点。

2.2 Smart scheduling：三条流水

把 $W$ 张卡切成 $W/\text{pipeline\_gran}$ 个 micro-batch，两条 CUDA stream 用 event 同步，形成三段重叠：

stream 0 (通信)：接收第 s+1 组的输入 token
stream 1 (计算)：对第 s   组已到的 token 跑 GEMM
stream 0 (通信)：把第 s-1 组算完的输出送回源 worker

2.3 Topology-aware gate：给跨节点流量设预算

节点内 NVLink 带宽比跨节点网卡高一个数量级。HirGate 直接给跨节点的出向 token 数设硬预算：

$\text{ogn\_threshold} = \lfloor S \times \text{ogn\_ratio} \rfloor \quad (\text{例如 batch 的} 14\%)$

正常算 Top-2 分数；一旦路由到远端节点的 token 数超过阈值，就把远端专家的分数置 $-\infty$ ，逼这些 token 改选节点内专家。只有 router logit 最高的那 $\text{ogn\_threshold}$ 个 token 保留跨节点的权限。

2.4 多层 roofline

把传统的算术强度 roofline 扩展成分布式多层天花板：

$\text{可达性能} = \min\big(\text{峰值算力},\ \text{AI}_{\text{DRAM}} \times \text{BW}_{\text{DRAM}},\ \text{AI}_{\text{Net}} \times \text{BW}_{\text{Net}}\big)$

其中 $\text{AI}_{\text{Net}} = \dfrac{\text{专家 FLOPs}}{\text{跨互连传输字节数}}$ 。分别对节点内 NVLink 和节点间 IB/以太网画天花板，就能定位通信延迟超过计算延迟的确切边界。

2.5 结果与遗留

64 卡（8 节点 × 8 V100/A100）上相对 PyTorch DDP / DeepSpeed ZeRO-2 / GShard / BASE Layer 加速 1.37×–17.87×；仅 dynamic shadowing 一项在强倾斜下就有 1.95×。

遗留问题：影子权重要在每张卡上多分配显存存权重和梯度；路由倾斜若在步间快速漂移，频繁广播会造成显存抖动。而且它仍在传统 batched GEMM + capacity factor 的框架内。

三、Tutel (MLSys 2023)

Hwang et al., arXiv:2206.03382

3.1 参数 $r$ ：把 DP/EP/MP 统一成一个连续旋钮

关键设计是权重布局不变：专家权重按 ZeRO-3 永久地切分在全部 $W$ 张卡上，每张卡持有唯一一片。切换并行策略时不搬运任何张量。

7 种 DP/EP/MP 组合里，Tutel 证明只有两种是可能最优的：Switchable DP（ $r=0$ ）和 Switchable EP+DP+MP（ $r \ge 1$ ）。参数 $r \in [0, \lceil W/E \rceil]$ 把 $W$ 张卡分组：

$\text{组大小} = \left\lceil \frac{W/E}{r} \right\rceil$

组内做 DP：前向 all-gather 权重，反向 reduce-scatter 梯度
组间做 MP：本地 token 复制 $r$ 份，各自算切片后的专家 GEMM，再跨 MP rank 求和

三个端点：

$r$	含义
$0$	纯 DP。权重切在全部 $W$ 卡上，token 全本地，无 all-to-all
$1$	标准 EP+DP。组大小 $W/E$
$\lceil W/E \rceil$	纯 EP+MP。组大小为 1，完全消除组内 DP 的权重 all-gather

这就是”激活通信 vs 参数通信”的旋钮： $r$ 增大，组内 all-gather 的参数通信减少，但 MP 的 token 复制和组间 all-to-all 的激活通信增加。

通信延迟关于 $r$ 在 $[1, \lceil W/E \rceil - 1]$ 上是凸的，所以用三分搜索，只要 $\log_{3/2}(\lceil W/E \rceil) + 2$ 次试探。

3.2 2DH All-to-All：四阶段

扁平 all-to-all 在千卡规模严重退化。2DH 拆成四步：

GPU 上做 stride 内存拷贝，把目标为同一张节点内 GPU 的非连续块打包成连续缓冲
节点内 all-to-all（NVLink，走 MSCCL/LL128 协议）
再一次 stride 拷贝，把目标为同一个远端节点的块打包
节点间 all-to-all，通信端点数从 $n$ 降到 $n/m$

3.3 自适应流水与动态容量

自适应流水只切 dispatch all-to-all、专家 GEMM、combine all-to-all 这三段，沿容量维 $C$ 切成 $d \in \{1,2,4,8\}$ 份；门控和残差连接不切，避免放大路由倾斜。

动态容量由 capacity_setting = x 控制：

$x > 0$ ：固定 capacity factor，丢 token
$x = 0$ ：零丢弃。每步检查路由结果，把容量设成恰好容纳最重专家所需： $c_{\text{step}} = \max_e (\text{分配给专家 } e \text{ 的 token 数})$
$x < 0$ ：动态容量但有硬上界 $|x|$

因为 $x=0$ 时 $c$ 每步都在变，Tutel 维护一个在线缓存 $\lfloor c/R \rfloor \to \{r^*, d^*, a^*\}$ （窗口 $R=128$ ），直接查出预搜好的最优并行度、流水度和集合算法，不必重新 profile。

3.4 结果与遗留

Azure 上从 16 卡扩到 2048 张 A100，单个 MoE 层相对 Fairseq/DeepSpeed 加速 4.96×（16 卡）、3.11×（128 卡）、5.75×（2048 卡）。真实模型 SwinV2-MoE 训练加速 1.55×、推理 2.11×。

遗留问题——这是最关键的一条： $x=0$ 的零丢弃是把所有专家的容量都设成 $\max_e T_e$ 。由于底层是 cublasGemmBatched 标准批量 GEMM，欠载的专家必须补零补到最大值，浪费 2×–11× 的算力，激活显存同样膨胀。

这条遗留问题直接催生了 MegaBlocks。

四、MegaBlocks (MLSys 2023)

Gale et al., arXiv:2211.15841

4.1 把 MoE 写成块稀疏矩阵乘

观察：专家 token 天然构成一个变尺寸的块对角激活矩阵。于是两层线性变成两个稀疏原语：

第一层 SDD（Sampled Dense-Dense）：稠密输入 $X$ 乘稠密权重 $W_1$ ，产出块对角的块稀疏中间激活 $H$ ：

$H = \text{SDD}(X, W_1, \text{topology})$

块 $(e,e)$ 的形状是 $(T_e \times d_{\text{ffn}})$ ， $T_e$ 是路由到专家 $e$ 的精确 token 数。

第二层 DSD（Sparse-Dense）：稀疏的 $H$ 乘稠密的 $W_2$ 得稠密输出：

$Y = \text{DSD}(H, W_2)$

反向用 $\text{SDD}^T$ 、 $\text{DSD}^T$ 、 $\text{DTSD}$ 、 $\text{DD}^T\text{S}$ 四个原语。

4.2 为什么是 BCSR + BCOO 混合

标准 BCSR 对 DSD 和 $\text{DD}^T\text{S}$ 的按行迭代很合适，但 SDD 并行执行时每个 threadblock 需要同时知道自己负责的输出块的行和列坐标。在 BCSR 的行偏移里搜索会带来线程发散和额外开销。

于是 MegaBlocks 混合编码：在 BCSR 的列索引之外，为每个非零块显式存一份行索引，让 threadblock 做 $O(1)$ 的坐标查表。

还有一个巧思——transpose indices：按转置序迭代 BCSR 通常要把所有非零值搬一遍内存；MegaBlocks 改为构造一个整数数组存转置序下非零块的偏移，用一层间接寻址实现零数据拷贝的转置迭代。

4.3 为什么块是 $128 \times 128$

在 A100 上跨 CUTLASS 各种 tile 配置实测， $128 \times 128$ （16384 个非零元）稳定达到峰值 Tensor Core 算术强度（与 cuBLAS 稠密 GEMM 持平），同时把稀疏元数据开销摊到可忽略（< 0.01% 显存）。

4.4 dMoE：padding 上界

对比 Tutel 的 $x=0$ ：容量取 $\max_e T_e$ ，浪费 $E \times \max_e T_e - T$ 个 token 的算力。

MegaBlocks 只需补到块边界：

$\text{每专家 padding} = (128 - (T_e \bmod 128)) \bmod 128 \le 127$

$E=64$ 时全系统 padding 浪费严格界在 $64 \times 127 = 8128$ 个 token，与路由倾斜无关。专家处理数千 token 时浪费 < 1%–3%，对比 Tutel 的 200%–1000%。

4.5 结果与遗留

8×A100 上训 64 专家 MoE，相对 Tutel 的补零式 dropless 端到端加速 1.38× / 2.0× / 4.35×（三个模型规模）；因为省下补零的激活显存，micro-batch 能开大 2× / 4× / 8×。18 个前反向问题形状上平均达到 cuBLAS 稠密吞吐的 98.6%。

遗留问题：仍要在 HBM 里显式分配并 materialize 置换补齐后的 token 缓冲和块稀疏中间激活矩阵。另外小 batch 或 $T_e < 128$ 时， $128\times128$ 的块量化会造成严重的边角利用率不足。

五、ScatterMoE (2024)

Tan et al., arXiv:2403.08245

5.1 MegaBlocks 到底做了几趟

先把被优化的对象说清楚。MegaBlocks 一层要跑五个阶段：

router 算 top- $k$
gather kernel 在 HBM 新分配一个补齐张量，把 token 按专家序拷进去
SDD 算 GEMM，把块稀疏激活写回 HBM
DSD 读稀疏激活，把置换序的输出写回 HBM
scatter kernel 把 token 反置换回序列序

5.2 ParallelLinear：四种模式

ScatterMoE 的核心是一个 Triton kernel _scatter2scatter，支持四种输入/输出组合：

模式	输入	输出
Scattered → Scattered	序列序	序列序
Scattered → Grouped	序列序	按专家连续
Grouped → Scattered	按专家连续	序列序
Grouped → Grouped	按专家连续	按专家连续

MoE MLP 的用法：第一层配成 Scattered → Grouped（中间态 $H$ 按专家连续存好，正合第二层胃口），第二层配成 Grouped → Scattered（直接写成最终的未置换序列布局，并顺手融合门控权重的乘加）。

于是三个缓冲全部消失：不分配置换输入、不 materialize 块稀疏中间矩阵和它的元数据、不分配中间的未置换输出。

5.3 routing index 怎么驱动寻址

预处理只排索引，不排数据：

flattened_expert_idxs = expert_idxs.flatten()
sorted_expert_idxs, sorted_scattered_idxs = torch.sort(flattened_expert_idxs)
expert_counts  = bincount(flattened_expert_idxs, minlength=num_experts)
expert_offsets = expert_counts.cumsum(-1)

kernel 内部（grid 在 $M$ = tokens × $k$ 上按 BLOCK_M=128 切）：

E_idxs = load(expert_idxs_ptr + M_block)     # 这批槽位属于哪个专家
M_idx  = load(grouped_idx_ptr + M_block)     # 排序后槽位 → 原 token 的置换

输入是 scattered 时的行号与指针：

$M_{\text{in\_idx}} = M_{\text{idx}} // k$ $X_{\text{ptr}} + M_{\text{in\_idx}}[:,\text{None}]\cdot \text{stride}_{xm} + K_{\text{block}}[\text{None},:]\cdot \text{stride}_{xk}$

这一步就是从 HBM 直接间接 gather 进 SRAM，没有独立的 gather kernel。

权重按专家 ID 索引：

$W_{\text{ptr}} + K_{\text{block}}[:,\text{None}]\cdot\text{stride}_{wk} + N_{\text{block}}[\text{None},:]\cdot\text{stride}_{wn} + E_{\text{idx}}\cdot\text{stride}_{we}$

输出是 scattered 时直接散写：

$Y_{\text{ptr}} + M_{\text{idx}}[:,\text{None}]\cdot\text{stride}_{ym} + N_{\text{block}}[\text{None},:]\cdot\text{stride}_{yn}$

累加器 acc 从寄存器直接散射到 HBM，不经中转。

5.4 结果与遗留

训练用 MegaBlocks 66.2% 的显存（降 33.8%），推理 53.6%（降 46.4%）。1.5B Mixtral 上因为能开 2× 的物理 micro-batch，训练吞吐高 38.1%；MoMHA 推理吞吐高 24.0%。

遗留问题：单机。不提供多机通信原语，分布式靠外部的 FSDP / Megatron。另外没有单 token 自回归解码的专用 kernel。

六、CCFuser (PPoPP 2025)

Wang, Xia, Yang, Zhou, Cheng. DOI:10.1145/3710848.3710868。代码 WongHuLin/CCFuser，无 arXiv 预印本。

6.1 用 NVSHMEM 而非集合通信

内存抽象是 NVSHMEM PGAS（2.9.0，跑在 MPI 上），对称堆用 nvshmem_malloc 分配。节点内走 NVLink/PCIe，跨节点可配 UCX/IB。

6.2 五个阶段，不是一个大 kernel

1. _local_scatter_kernel      按门控把 token 重排进本地对称缓冲
2. nvshmemx_barrier_all_on_stream
3. moe_first_mm               融合 dispatch + GEMM1
4. 写回全局内存 + PyTorch 激活函数
5. moe_second_mm              融合 GEMM2 + combine
6. barrier + index_copy_      反置换回序列序

6.3 dispatch 用 pull，combine 用 push

dispatch（GEMM1 内）：专家 CTA 发一个非阻塞的设备端读 nvshmemx_get16_nbi_block，把远端 token 拉进本地 merged_input。传输在飞的同时，同一个 CTA 先算本地那 64 个 token 的 GEMM1，然后 nvshmem_quiet() 等远端数据到齐再算远端那块。

combine（GEMM2 内）：CTA 先算远端 token 的 GEMM2 进临时缓冲 temp_remote_mem，发非阻塞写 nvshmemx_put128_nbi_block 推回目标 rank，然后再算本地 token。

CTA 内是 256 线程 8 warp 的同构协作——没有 warp 特化，所有 warp 一起发传输、一起算 Ampere WMMA。tile 是 $BM=64, BN=512, BK=32$ 。

6.4 结果与遗留

A100 集群（NVLink 与 PCIe 两种配置），分布式训练（含反向与权重梯度 kernel），EP=4 为主、脚本支持到 EP=8。相对 FastMoE 平均 2.96×、FasterMoE 平均 2.48×，峰值 4.34×。

四条局限（作者自述）：

固定 $BM=64$ 离散化：token 数不是 64 的倍数就要补齐，强倾斜下碎片明显
本地/远程块配对依赖：靠”一个本地块配一个远程块”在 CTA 内互相掩藏延迟；某张卡本地块过少就无从掩藏，需要离线的 expert reassignment 图变换来再平衡
中间激活全量落 HBM：merged_input、ffn1_out 用 torch::zeros 分配，供 autograd 保存——这是做训练的必然代价
主机在关键路径上：多次 kernel launch 加 stream barrier

七、SonicMoE (2025.12)

Guo, Mishra, Cheng, Stoica, Dao. arXiv:2512.14080

7.1 问题：细粒度 MoE 的两个标度

激活显存：ScatterMoE 和 MoMoE 为反向缓存 $O(TKd)$ 的中间张量，专家粒度越细，每层激活显存线性增长。

IO：细粒度稀疏 MoE 在访存受限区。论文的说法是「增大 $G$ 和增大稀疏度（减小 $\rho$ ）都会把算术强度往下压」，于是 kernel 运行时间由 IO 主导而非算力。

7.2 唯一改了数学的一条：反向收缩重排

教科书做法是先得到前向输出 $Y$ 再点积算门控梯度， $Y$ 必须缓存。SonicMoE 重排收缩顺序：

$dS_{t,e} = \langle dO_t,\ Y_{e,t}\rangle = \langle dA'_{e,t},\ A_{e,t}\rangle$

用两个已有的中间量直接收缩出来， $Y$ 从不落地。效果是激活显存不随粒度 $G$ 增长。

7.3 Blackwell 上的四件工程

warp 特化：一个 producer warp 发 gather；一个 MMA warp 异步发 tcgen05.mma 进 TMEM；多个 epilogue warp 用 tcgen05.ld 在 MMA 继续跑的同时排空 TMEM。

TMEM 管理：每 SM 128 行 × 512 列 × 32bit = 256KB，512 列正好放两个各 256 列的独立累加阶段。

2CTA MMA + cp.async 的同步难题：cp.async 只能在本 CTA 内发信号，但 leader CTA 的 MMA 需要两个 CTA 的数据都就绪。解法是 relay warp——CTA 1 里 1 个 warp 取索引、4 个 warp 发 cp.async gather、1 个 relay warp 等本 CTA 的 cp.async 完成后去 CTA 0 的 barrier 上 arrive。

gather 融合双路径：cp.async 与 TMA gather4（cp.async.bulk.tensor.2d.shared::cta.global.tile::gather4）两条路，运行时 autotune 选。并且 gather + SwiGLU 一起塞进 up-projection 的 epilogue。

CLC 调度：用 clusterlaunchcontrol.try_cancel 让运行中的 cluster 向硬件问下一个 tile 坐标，不碰 GMEM atomic。

7.4 结果

B300 上，前向/反向 TFLOPS 比 DeepGEMM 基线高 54%/35%，比 Triton 官方样例前向高 21%，相对 ScatterMoE 和 MoMoE 吞吐翻倍（端到端 1.87–4.04×）。激活显存在粒度增大时保持恒定。

八、当前落点：DeepGEMM Mega MoE (2026.04)

放在这条线上看，它的位置是：把 ScatterMoE 遗留的两条（无多机原语、无解码 kernel）当作起点，代价是在单卡内融合得比 ScatterMoE 更浅。

细节见另一篇，这里只列它在本文脉络里的四个选择：

通信：全 pull，用 PyTorch symmetric memory 的对等指针（纯 NVLink，kNumMaxRanks = 72），不做集合通信
中间激活：有界环形缓冲，容量按「同时在飞的 block 数」定而非按路由后总 token 数
tile 粒度：候选 {8,16,32,64,96,128,192}，按 $\bar{T}_e$ 分档，最小到 16——正面回应 MegaBlocks 承认的小 batch 短板
调度：全局原子计数器工作队列，因为任务总数要等 kernel 内的跨卡计数交换才知道，用不了 CLC

九、收成八条设计决策

决策 0：先算 $\bar{T}_e$

$\bar{T}_e = \frac{\text{num\_tokens} \times \text{num\_ranks} \times \text{num\_topk}}{\text{num\_experts}}$

$\bar{T}_e < 128$ 时 MegaBlocks/ScatterMoE 的固定 128 tile 有严重 fringe； $\bar{T}_e$ 很大时小 tile 又损失 Tensor Core 效率。这个数把设计空间劈成两半。

决策 1：容量策略

做法	出处	代价
静态 capacity factor 丢 token	GShard / DeepSpeed-MoE / FasterMoE	掉点
补零到 $\max_e T_e$	Tutel $x=0$	浪费 2–11× 算力
块稀疏补到块边界	MegaBlocks	≤127 token/专家；要写 SDD/DSD
指针间接寻址零 padding	ScatterMoE	只 pad 索引数组
最坏情况预留池	Mega MoE	显存换确定性

决策 2：tile 粒度

固定 128（训练、 $\bar{T}_e \gg 128$ ）或按 regime 选档（推理）。后者可抄 DeepGEMM 的表： $\bar{T}_e \le 8.5 \to 16$ 、 $\le 16.5 \to 32$ 、 $\le 32.5 \to 64$ 、 $\le 64.5 \to 96$ 、 $\le 96.5 \to 128$ 、其余 $192$ 。最小档配大 BLOCK_K（256）摊薄同步开销。

决策 3：token 搬运方向

判据是有没有对等内存。 没有 → 集合通信，优化空间限于分层（DeepSpeed-MoE）和 2DH（Tutel）。有 → pull（对等指针最轻但锁死 scale-up；NVSHMEM 能出节点但要引运行时）。

配套：rank 轮转交错（min-peeling）

走 pull 的话，本地某个专家的 token 池别按 rank 拼接。设 $R$ 个 rank 各发来 $c_r$ 个 token，每轮取活跃 rank 数 $A$ 与剩余量最小值 $L$ ，发出 $L\times A$ 个槽位在活跃 rank 间轮转，然后各减 $L$ （至少一个 rank 归零），最多 $R$ 轮结束。

4 个 rank、 $[5,3,3,1]$ 的例子：

轮	活跃 $A$	最小 $L$	发出	剩余
1	4	1	4 槽	$[4,2,2,0]$
2	3	2	6 槽	$[2,0,0,0]$
3	1	2	2 槽	$[0,0,0,0]$

槽位:  0   1   2   3  |  4   5   6   7   8   9  | 10  11
rank: r0  r1  r2  r3  | r0  r1  r2  r0  r1  r2  | r0  r0
序号:  0   0   0   0  |  1   1   1   2   2   2  |  3   4

给定槽位反查的实现（impls/sm100_fp8_fp4_mega_moe.cuh:461）：

const uint32_t slot_idx_in_round = slot_idx % num_active_ranks;   // 选哪个 rank
token_idx_in_rank = offset + (slot_idx / num_active_ranks);        // 选它的第几个 token

活跃集合每轮在变，所以选 rank 那步要用 __ballot_sync 做 mask、__fns 找第 $n$ 个置位，不能直接下标。

为什么值得：两种排法从每个 rank 读的总字节数完全一样，差别在同一时刻在读谁。dispatch 循环的步长是全卡的 dispatch warp 总数：

constexpr uint32_t kNumGlobalWarps = kNumSMs * kNumDispatchWarps;
for (uint32_t token_idx = sm_idx * kNumDispatchWarps + warp_idx; ;
     token_idx += kNumGlobalWarps)

所以任意时刻整卡在处理一个约 kNumGlobalWarps 宽的连续槽位窗口（132 SM × 4 warp ≈ 500 槽同时在飞）。按 rank 拼接时这 500 个槽很可能全落在同一个 rank 的区段里——500 个并发 pull 打向同一块远端 GPU，其余链路闲置，链路被串行使用；轮转交错则让 8 条 NVLink 同时在跑。

买到的是瞬时链路并发度，代价只是一段 warp 内位运算。

两个边界：交错的只是池子布局，源元数据仍按 (专家, rank, rank 内序号) 存；它不解决专家之间的计算不均（那是决策 7），对 GEMM 侧完全中性。

决策 4：中间激活

由「要不要反向」决定，没有自由度。 要反向 → 落盘（CCFuser 的选择是被 autograd 逼的），或用收缩重排从算法上绕开（SonicMoE）。只做前向 → 不 materialize，或有界环。

决策 5：融合层次

层次	出处	失去
跨 stream 编排	FasterMoE / Tutel	CPU 在关键路径；kernel 边界断流水
分阶段跨 kernel	CCFuser	多次 launch；中间激活必须落盘
单 kernel warp 特化	Mega MoE / SonicMoE	反向难做；寄存器预算紧

走最后一条必须加寄存器静态断言，上限 64512/SM，否则 occupancy 静默掉档。

决策 6：调度机制

判据一句话：路由影响的是寻址，还是 grid 本身？

路由影响寻址   → 静态 grid + 间接寻址就够   → 能用 CLC      （SonicMoE / ScatterMoE）
路由影响 grid  → 只能运行时发现工作量        → 只能原子队列  （Mega MoE）

为什么单卡算子的 tile 数是静态的

单卡 grouped GEMM 的三个维度全是纯 shape：

$M = L \times K \ (\text{token 数} \times \text{top-k}), \qquad N = d_{ff}, \qquad K = d$

关键是 $M = L\times K$ 这个恒等式：每个 token 恰好产生 $K$ 个 (token, expert) 任务槽，一个不多一个不少。路由决定的是每个槽用哪个专家的权重、读原张量的哪一行——那是 tile 内部的间接寻址：

E_idxs = load(expert_idxs_ptr + M_block)   # 用哪个专家的权重
M_idx  = load(grouped_idx_ptr + M_block)   # 读哪些行

唯一与路由有关的 tile 数变化，是按专家补齐到块边界带来的，上界为 $E$ 个 tile。所以 grid 按 $\lceil LK/\text{BLOCK\_M}\rceil + E$ 静态发即可，多出的几块早退。路由在 kernel 之外就算完了。

为什么融了通信之后就不行

Mega MoE 把跨卡 dispatch 融进 kernel，于是「本地专家会收到多少 token」取决于别的 GPU 路由了什么，而那需要一次跨卡交换——那次交换现在在 kernel 内部：

do {
    value = ptx::ld_volatile(workspace.get_expert_recv_count_sum_ptr(expert_idx));
} while (static_cast<uint32_t>(value >> 32) != kNumSMs * kNumRanks);
...
num_total_m_blocks = get_num_total_pool_blocks();

等的是所有 rank 的所有 SM 都报完数。在这行之前，块数不存在。

按最坏情况发 grid 理论上可行，但超发倍数是

$\frac{T_{\max}\cdot N_{\text{ranks}}\cdot\min(k, E_{\text{local}})}{T\cdot k} \ \approx\ N_{\text{ranks}} \quad (T = T_{\max},\ k \le E_{\text{local}})$

EP8 是 8×，EP32 是 32×，再叠 decode 时 $T_{\max}$ 相对实际 $T$ 的容量余量。

三条结构性障碍（即便愿意吃超发也绕不开）

阶段屏障。所有 CTA 必须先做完 dispatch 才能取 GEMM 任务（grid_sync 于 :380，nvlink_barrier 于 :403）。CLC 的模型里 CTA 是可互换的 tile worker，没有「阶段」概念。
两条互相依赖的任务流。L2 任务要等对应 L1 任务先被取走：while (ld_volatile(l1_task_count) < num_required_l1_tasks) {}。CLC 只给一条线性化 tile 流。
多个 phase 各有独立计数器，还要控制发放节奏（warm-up wave，get_num_l1_warmup_waves）。

一般化的结论：决定落在哪一边的，是路由的最终解析发生在 kernel 边界的哪一侧。融合通信等于把「路由的最后一公里」拖进 kernel——从那一刻起，CLC、静态 tile 网格、host 侧的一切形状推断全部失效。这是融合通信的隐性税。

决策 7：负载不均在哪吸收

位置	出处	代价
参数层：复制热门专家	FasterMoE dynamic shadowing	额外显存；要在线成本模型
门控层：给跨节点流量设预算	FasterMoE HirGate	改变路由结果，可能影响精度
编译期：离线专家重分配	CCFuser	需预知负载分布
数据布局层：rank 轮转交错	Mega MoE	零成本，但只摊平带宽
路由层：逻辑→物理冗余映射	EPLB	需要额外物理副本

实务顺序：先布局层（零成本）→ 再路由层 → 最后才复制专家。

决策 8：scale-up 还是 scale-out

EP 度数会不会超出 NVLink 域？不会 → 对等指针最轻。会 → NVSHMEM，或域内对等 + 跨域 RDMA 分层。

反模式

拿 $\bar{T}_e$ 很小的场景测固定 128 tile 的 kernel。 MegaBlocks 自己承认 $T_e<128$ 时 fringe 严重，测错了目标。
把「dropless」当单一概念。 至少三种：Tutel 的补零到 $\max_e T_e$ 、MegaBlocks 的块稀疏、Mega MoE 的最坏情况预留，代价完全不同。
在需要反向的路径上追求「不 materialize」。 autograd 要存激活，要么落盘要么用重算/收缩重排绕开，没有第三条路。
假设融合越深越好。 Mega MoE 在单卡内比 ScatterMoE 融得浅——token 进两趟全局显存——它把预算花在跨卡上了。融合有预算，要花在瓶颈那一侧。
忽略 DeepSpeed-MoE 那个前提。 如果你的部署没把专家分散到足够多卡上聚合带宽，后面所有 kernel 优化都是在优化一个仍然受 HBM 带宽约束的系统。