GC-MoE
An MoE inference operator that decouples communication and computation using a resource-partitioning mechanism based on CUDA Green Contexts. By dividing the GPU into separate communication and computation partitions, it improves the efficiency of expert-parallel decoding.
Timeline (last 90 days)
GC-MoE is an open-source FP8/FP4 MoE inference operator for NVIDIA Hopper and Blackwell GPUs, which uses CUDA Green Contexts to partition the device into two partitions: a small communication partition and a large compute partition.
On the B300, GC-MoE allocates only 20 of the 148 SMs to dispatch and leaves all the rest for computation.
On 8 H20 GPUs, GC-MoE accelerates DeepSeek-V4.1-Flash decode by an average of 1.188x across the routed and shared-expert suites, with peaks of 1.296x and 1.353x respectively.
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 the 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 complete; the double-buffered receive counters eliminate the tail cross-rank barrier between replays.
The GC-MoE compute partition runs operand-swapped tcgen05 block-scaled MMA on Blackwell, and on Hopper runs register-A WGMMA with FP4 decoding completed in registers; both fuse SwiGLU, requantization, and optional shared experts.
The GC-MoE approach depends on Hopper/Blackwell-generation hardware features, and the partition ratio (such as 20/148 on the B300) needs to be tuned for specific models and topologies; for inference services with very small decode batches, the communication partition share is a key parameter that requires empirical parameter sweeping.
GC-MoE offers an engineering paradigm: use hardware-level resource partitioning to thoroughly isolate different responsibilities so that their respective scheduling policies do not interfere with each other, rather than cramming all stages into one kernel and overlapping them by sharing SMs.
All Facts (8)
GC-MoE is an open-source FP8/FP4 MoE inference operator for NVIDIA Hopper and Blackwell GPUs, which uses CUDA Green Contexts to partition the device into two partitions: a small communication partition and a large compute partition.
50%UnverifiedOn the B300, GC-MoE allocates only 20 of the 148 SMs to dispatch and leaves all the rest for computation.
50%UnverifiedOn 8 H20 GPUs, GC-MoE accelerates DeepSeek-V4.1-Flash decode by an average of 1.188x across the routed and shared-expert suites, with peaks of 1.296x and 1.353x respectively.
50%UnverifiedAt 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 the fused shared experts).
50%UnverifiedThe 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 complete; the double-buffered receive counters eliminate the tail cross-rank barrier between replays.
50%UnverifiedThe GC-MoE compute partition runs operand-swapped tcgen05 block-scaled MMA on Blackwell, and on Hopper runs register-A WGMMA with FP4 decoding completed in registers; both fuse SwiGLU, requantization, and optional shared experts.
50%UnverifiedThe GC-MoE approach depends on Hopper/Blackwell-generation hardware features, and the partition ratio (such as 20/148 on the B300) needs to be tuned for specific models and topologies; for inference services with very small decode batches, the communication partition share is a key parameter that requires empirical parameter sweeping.
50%UnverifiedGC-MoE offers an engineering paradigm: use hardware-level resource partitioning to thoroughly isolate different responsibilities so that their respective scheduling policies do not interfere with each other, rather than cramming all stages into one kernel and overlapping them by sharing SMs.
50%