GC-MoE：通过 CUDA Green Contexts 重叠专家并行令牌分发与 FP8/FP4 分组 GEMM

在专家并行模式下，混合专家（MoE）层的每一层都需要执行全到全令牌分发、分组 GEMM 以及合并步骤。在解码阶段，每个 rank 仅持有 1–256 个令牌：分发延迟和权重数据传输占据主导，而张量核心利用率本身就较低。在 DeepGEMM 的 Mega MoE 等融合算子中，分发与计算只有通过共享每个 SM 才能实现重叠：分发 warp 会与 TMA/MMA 和 epilogue warp 争用寄存器、共享内存及指令发射槽；每个阶段共享同一份 CTA 资源预算，同步范围覆盖整个 SM 阵列，并且每次重放都以工作区清理遍和跨 rank 屏障结束。我们提出 GC-MoE，这是一种面向 NVIDIA Hopper 和 Blackwell GPU 的开源 FP8/FP4 MoE 推理算子，利用 CUDA Green Contexts 将设备划分为一个小型通信分区和一个大型计算分区。每种角色都拥有独立的内核、资源预算和同步域，因此分发与 GEMM 可以分别进行调优。一个包含三个节点的 CUDA Graph（在启用融合共享专家的 Hopper 上为四个节点）将分发内核和分组 GEMM 内核作为并发根节点启动；令牌环握手使计算从首批已填充的环块开始，双缓冲接收计数器则消除了重放之间末尾的跨 rank 屏障。计算分区在 Blackwell 上运行操作数交换的 tcgen05 块缩放 MMA，在 Hopper 上运行采用寄存器 A 的 WGMMA，并在寄存器中执行 FP4 解码，同时融合 SwiGLU、重新量化和可选的共享专家。在 8 张 H20 GPU 上，GC-MoE 使 DeepSeek-V4.1-Flash 解码在路由专家和共享专家测试集上的平均速度均提升 1.188 倍（峰值分别为 1.296 倍和 1.353 倍）；在 8 张 B300 GPU 上，它使 GLM-5.2 和 DeepSeek-V4.1-Flash 的速度提升 1.106 倍（融合共享专家时为 1.123 倍，峰值为 1.209 倍），在准确率匹配的条件下，在全部 54 个测量点上均胜过 Mega MoE。GC-MoE 在 B300 上仅为分发保留 148 个 SM 中的 20 个。