为什么值得读
MoE decode 阶段的算子设计有个结构性矛盾：每 rank 只有 1–256 个 token，分发延迟和专家权重流量主导耗时，而张量核利用率天然低。DeepGEMM 的 Mega MoE 走「全 SM 融合」路线，代价是 dispatch warp 与 TMA/MMA/epilogue warp 争抢寄存器、共享内存和发射槽，同步覆盖整个 SM 阵列。GC-MoE 用 Green Contexts 把两个角色物理隔离、独立调优，这是个值得记住的工程取舍：牺牲 13.5% 的 SM，换来同步域解耦。Blackwell tcgen05 与 Hopper WGMMA 双后端都做了，含寄存器内 FP4 解码这类细节。
问题：decode 阶段 MoE 层的三笔开销与融合路线的代价
专家并行下每层 MoE 要付 all-to-all token 分发、分组 GEMM 与 combine 三笔开销。decode 时每 rank 仅 1–256 token，分发延迟与权重流量主导，张量核利用率天然低。
DeepGEMM Mega MoE 这类融合算子靠「共享每个 SM」做重叠，dispatch warp 与 TMA/MMA、epilogue warp 争抢寄存器/共享内存/发射槽，一次同步横跨整个 SM 阵列，且每次重放以 workspace 清理与跨 rank 屏障收尾。
方法：分区隔离 + 三节点 CUDA Graph
GC-MoE 用 CUDA Green Contexts 把设备拆成小通信分区与大计算分区，每个角色有自己的内核、资源预算与同步域，分发与 GEMM 可独立调优。
一个三节点 CUDA Graph（Hopper 上含融合共享专家时为四节点）把分发与分组 GEMM 内核作为并发根启动；token-ring 握手让计算在第一批填满的 ring block 上即启动；双缓冲接收计数器消除重放之间的跨 rank 尾屏障。
计算分区在 Blackwell 上跑 operand-swapped tcgen05 block-scaled MMA，在 Hopper 上跑寄存器内 FP4 解码的 WGMMA，融合 SwiGLU、重量化与可选共享专家。
结果：两个 GPU 世代、两个模型上的一致收益
8×H20：DeepSeek-V4.1-Flash decode 平均 1.188×（routed 与 shared-expert 两套套件同值，峰值分别 1.296× 与 1.353×）。
8×B300：GLM-5.2 与 DeepSeek-V4.1-Flash 平均 1.106×（融合共享专家后 1.123×，峰值 1.209×），在匹配精度下 54 个测量点全部超过 Mega MoE。通信分区仅占 B300 148 个 SM 中的 20 个。
边界
基准模型为 DeepSeek-V4.1-Flash 与 GLM-5.2，其他 MoE 架构的收益未在摘要中给出；对比对象限于 DeepGEMM Mega MoE。
摘要未给出端到端服务级指标（如 token 吞吐/延迟分布），收益限于算子层面。
本站仅据 arXiv 摘要页（E1）整理，未复现。摘要称算子为开源，但摘要页未附代码仓库链接（截至快照日），读者请以 PDF 全文与作者发布渠道为准；对比基线仅 DeepGEMM Mega MoE，端到端服务指标未提供。
本文基于论文原文摘要与方法章节整理，观点由编辑撰写；如与原文有出入，以 arXiv 原文为准。 内容为信息整理与技术评论，不构成投资或决策建议。 发现错误请通过站内纠错页提交（记录可查），修订记录只增不删。