GC-MoE 用 CUDA Green Contexts 加速 MoE 推理，DeepSeek 与 GLM 解码提速超 10%
GC-MoE: Overlapping Expert-Parallel Token Dispatch with FP8/FP4 Grouped GEMM via CUDA Green Contexts
给搞 MoE 推理优化的人看的新算子，8 卡 H20 上跑 DeepSeek 解码快了近 20%，还开源了，可以直接对比 Mega MoE 试一下。
GC-MoE 是一个开源 FP8/FP4 MoE 推理算子，支持 NVIDIA Hopper 和 Blackwell GPU，利用 CUDA Green Contexts 把设备拆成通信与计算两个独立分区。在 8 张 H20 上，DeepSeek-V4.1-Flash 解码平均加速 1.188 倍，峰值达 1.353 倍。在 8 张 B300 上，GLM-5.2 和 DeepSeek-V4.1-Flash 平均加速 1.106 倍，在全部 54 个测量点上超过 DeepGEMM 的 Mega MoE。通信分区在 B300 上仅占用 148 个 SM 中的 20 个。
GC-MoE: Overlapping Expert-Parallel Token Dispatch with FP8/FP4 Grouped GEMM via CUDA Green Contexts
Under expert parallelism, mixture-of-experts (MoE) layers pay all-to-all token dispatch, grouped GEMMs, and a combine step in every layer. In decode each rank holds only 1-256 tokens: dispatch latency and weight traffic dominate while tensor-core utilization is inherently low. In fused operators such as DeepGEMM's Mega MoE, dispatch and computation overlap only by sharing every SM: dispatch warps compete with TMA/MMA and epilogue warps for registers, shared memory, and issue slots; every phase shares one CTA resource budget, synchronization spans the whole SM array, and every replay ends with a workspace-clean pass and a cross-rank barrier. We present GC-MoE, an open-source FP8/FP4 MoE inference operator for NVIDIA Hopper and Blackwell GPUs that uses CUDA Green Contexts to split the device into a small communication partition and a large compute partition. Each role gets its own kernel, resource budget, and synchronization domain, so dispatch and GEMM tune independently. A three-node CUDA Graph (four on Hopper with fused shared experts) launches the dispatch and grouped-GEMM kernels as concurrent roots; a token-ring handshake starts computation on the first filled ring blocks, and double-buffered receive counters remove the trailing cross-rank barrier between replays. The compute partition runs operand-swapped tcgen05 block-scaled MMA on Blackwell and register-A WGMMA with in-register FP4 decoding on Hopper, with fused SwiGLU, requantization, and optional shared experts. On 8 H20 GPUs GC-MoE accelerates DeepSeek-V4.1-Flash decode by 1.188x on average on both the routed and shared-expert suites (peaks 1.296x and 1.353x); on 8 B300 GPUs it accelerates GLM-5.2 and DeepSeek-V4.1-Flash by 1.106x (1.123x with fused shared experts, peak 1.209x), beating Mega MoE at all 54 measured points at matched accuracy. GC-MoE reserves only 20 of 148 SMs on B300 for dispatch.
