GC-MoE: Decoupling MoE Communication and Computation with CUDA Green Contexts

Why Expert Parallelism Is Inefficient in MoE Decoding

Under expert parallelism, every MoE layer goes through three stages: all-to-all token dispatch, grouped GEMM, and combine. In the decode phase, each rank holds only 1 to 256 tokens, so per-GPU computation produces very little, while weight movement and communication latency stay exactly the same.

This means MoE in the decode scenario has a classic "communication and computation dragging each other down" structure. Fused operators like DeepGEMM's Mega MoE attempt to hide latency through overlap, but their approach is to let communication warps and computation warps share the same pool of SMs: registers, shared memory, and issue slots all compete together, the synchronization domain spans the entire SM array, and every replay has to clean up the workspace and perform a cross-rank barrier. Overlap is achieved, but the resource budgets of the two sides cannot be tuned independently.

All-to-all communication is the core overhead of expert parallelism: each token must be sent, based on gating routing results, to the GPU holding the corresponding expert, and then collected after computation. In the decode phase, each GPU has only a handful of tokens per iteration; although the all-to-all payload is small, the fixed latency (kernel launch, NVLink/IB transfer handshakes) cannot be amortized, so communication's share of time rises sharply.

Grouped GEMM refers to batching the small matrix multiplications of multiple experts into a single kernel call to avoid the overhead of launching a kernel per expert. However, the number of tokens assigned to each expert is highly uneven, leading to load imbalance and reduced SM utilization.

Mega MoE / fused operators take the approach of cramming dispatch, GEMM, and combine into a single persistent kernel, using warp specialization to overlap communication and computation and eliminating kernel-launch and intermediate-workspace round trips—but at the cost of resource contention and coupled synchronization domains.

GC-MoE's Core Idea: Split the Device into Two Independent Partitions

GC-MoE is an open-source FP8/FP4 MoE inference operator for NVIDIA Hopper and Blackwell GPUs. It uses CUDA Green Contexts to partition the device into two parts: a small communication partition and a large computation partition.

The key benefit of this split: each role has its own kernel, its own resource budget, and its own synchronization domain. Dispatch and GEMM can be tuned independently without competing for resources. On B300, GC-MoE allocates only 20 of the 148 SMs to dispatch and leaves the rest entirely for computation.

Three-Node CUDA Graph and token-ring Handshake

At the scheduling level, GC-MoE uses a three-node CUDA Graph to launch the dispatch kernel and the grouped-GEMM kernel as concurrent root nodes (on Hopper, the graph has four nodes because of fused shared experts). The token-ring handshake lets computation begin as soon as the first batch of ring blocks is filled, rather than waiting for the entire dispatch to finish; double-buffered receive counters eliminate the trailing cross-rank barrier between replays.

CUDA Graph captures a series of kernel launches and their dependencies as a static graph; a single submission lets the driver schedule them by dependency, eliminating per-kernel CPU launch overhead and launch gaps. "Three nodes" means the graph contains three kernel nodes—dispatch, grouped-GEMM, and combine/finalization—where the first two are launched concurrently as root nodes, and the token-ring handshake establishes fine-grained producer-consumer synchronization on the device side rather than relying on graph-level serial edges.

token-ring handshake is a ring-buffer-style cross-kernel synchronization: each time dispatch fills a batch of ring blocks it increments a counter, and the compute kernel polls that counter, starting GEMM as soon as the first batch of data is ready, thereby pipelining computation and dispatch. Double-buffered receive counters let the next replay reuse the previous round's buffers without inserting a cross-rank barrier at the end of each round.

Two MMA Implementations in the Compute Partition

The compute partition takes different routes for the two hardware generations: on Blackwell it runs operand-swapped tcgen05 block-scaled MMA, and on Hopper it runs register-A WGMMA with in-register FP4 decoding. Both fuse SwiGLU, requantization, and optionally shared experts.

CUDA Green Contexts is a hardware-level resource partitioning mechanism introduced in CUDA 12.4 that allows a GPU's SMs to be divided into several non-overlapping sets, with each context having its own SM quota. Unlike MPS or stream priorities, Green Contexts partitioning is spatially isolated: kernels on different partitions do not compete for the same SMs, and their occupancy, shared memory, and register budgets do not affect each other. This lets communication kernels and compute kernels be tuned independently for their own optimal block size, occupancy, and synchronization strategy without conceding to the other's resource needs. This feature is currently fully supported only on Hopper and newer architectures.

WGMMA (Warpgroup Matrix Multiply-Accumulate) is the asynchronous tensor core instruction introduced in the Hopper architecture, executed at the granularity of a warpgroup of 128 threads, with operand A able to come from registers or shared memory. register-A means the A operand resides in registers, avoiding shared-memory round trips, which suits the small M dimension of the decode phase.

tcgen05 MMA is the fifth-generation tensor core instruction set of the Blackwell architecture, introducing Tensor Memory (TMEM) as a dedicated storage hierarchy for accumulators and operands, and supporting block-scaled formats—i.e., FP4/FP8 microscaling (MX) formats, where each 32-element block shares an exponent scaling factor, compressing weight bit-width to 4 bits while preserving accuracy. operand-swapped means swapping the roles of the A/B operands so that weights enter the MMA in a more favorable layout, reducing layout-conversion overhead.

Measured Results: Consistent Speedups and a Clean Sweep

On 8 H20 GPUs, GC-MoE speeds up DeepSeek^Ad-V4.1-Flash decode by an average of 1.188× across the routed and shared-expert suites, with peaks of 1.296× and 1.353× respectively. On 8 B300 GPUs, it delivers a 1.106× speedup on GLM-5.2 and DeepSeek-V4.1-Flash, 1.123× with fused shared experts, and a peak of 1.209×—and it beats Mega MoE at the same precision on all 54 measurement points.

These numbers may not look dramatic in magnitude, but considering that Mega MoE is itself a highly optimized fused operator, maintaining a consistent lead across all measurement points shows that the independent tuning space enabled by resource isolation is genuinely effective, rather than a lucky win under one particular configuration.

An Observation on MoE Inference Engineering

GC-MoE's value lies not just in its speedup but in the engineering paradigm it demonstrates: rather than cramming all stages into a single kernel and overlapping them by sharing SMs, it is better to use hardware-level resource partitioning to fully isolate different responsibilities so that their scheduling strategies do not interfere. Green Contexts were originally used mostly for resource reservation in latency-sensitive scenarios; here they are used to decouple communication and computation—an approach worth learning from.

Of course, this approach also depends on the Hopper/Blackwell hardware generation, and the partition ratio (e.g., 20/148 on B300) needs to be tuned for the specific model and topology. For inference services with very small decode batches, the communication partition's share is a key parameter that requires empirical sweep tuning.

One-line Summary

GC-MoE uses CUDA Green Contexts hardware partitioning to decouple communication and computation in MoE decoding, surpassing Mega MoE across the board.

Paragraph Summary

The article analyzes the root cause of low expert-parallelism efficiency in MoE decoding: each GPU has very few tokens, yet weight movement and all-to-all communication latency remain unchanged, while fused operators like Mega MoE make communication and computation share SMs so resources cannot be tuned independently. GC-MoE uses CUDA Green Contexts to divide Hopper/Blackwell GPUs into an independent communication partition and computation partition, each with its own kernel, resource budget, and synchronization domain. For scheduling, it uses a three-node CUDA Graph to launch dispatch and grouped-GEMM concurrently, achieving pipeline overlap via a token-ring handshake, with double-buffered counters eliminating the tail barrier. The compute partition uses tcgen05 block-scaled MMA on Blackwell and register-A WGMMA with in-register FP4 decoding on Hopper. Measured results show an average 1.188× speedup on H20 and 1.106× on B300, leading Mega MoE at the same precision across all 54 measurement points.