GC-MoE: Overlapping Expert-Parallel Token Dispatch with FP8/FP4 Grouped GEMM via CUDA Green Contexts

Yong Li; Kan Liu; Tao Lan

cs.DCNov 3, 2026

GC-MoE: Overlapping Expert-Parallel Token Dispatch with FP8/FP4 Grouped GEMM via CUDA Green Contexts

Authors: Yong Li, Kan Liu, Tao Lan

Abstract

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.

Figures & tables

Panel of Figure 1 — Figure 1: Execution timelines. (a) On B300 the baseline is DeepGEMM’s Mega MoE, one persistent fused kernel on all 148 SMs: after routing counts and a tag-1 cross-rank barrier, the dispatch warps stream tokens, scales, and routing metadata into a per-block token ring while the GEMM consumes each ring block as soon as it is full. The phases overlap, but on the same SMs, with the dispatch warps and their staging buffers carved out of every SM’s register file and shared memory; the tail is $\max(T_{\mathrm{combine}},t_{\mathrm{clean}}+t_{\mathrm{tag2}})$ —the workspace clean overlaps the combine, but the tag-2 barrier is still a cross-rank round trip on every replay. (b) GC-MoE runs role1 (dispatch) on the 20-SM communication context and role2 (grouped GEMM) plus role3 (combine) on the 128-SM compute context. Role2 starts when the first ring block is filled (dashed trigger), and parity-double-buffered receive counters remove the tag-2 barrier, so the next replay can start as soon as role3 completes (green band). Each role is an ordinary kernel with its own register, shared-memory, and synchronization budget, so the dispatch and GEMM paths can be optimized independently. Bars are illustrative; measured end-to-end speedups span 1.03–1.30 $\times$ across the routed-expert sweeps and 1.01–1.35 $\times$ including shared experts (Section 6.2 ). Partitions are the B300 values.

Role	Function	Green context	Grid	Threads	Cluster
role1	token dispatch (all-to-all pull)	communication	dispatch_sms	256 or 384	–
role2	FP8 $\times$ FP4 grouped GEMM + SwiGLU	compute	compute_sms	384	(2,1,1)
role3	signal reset + combine	compute	finalize_sms	384	–

Table 1: Roles of the SM103 CUDA Graph (three nodes; on SM90, fused shared experts add a fourth node). dispatch_sms / compute_sms are the driver-aligned partition sizes. finalize_sms is

\min(\texttt{compute\_sms},64)

at up to 128 tokens per rank and compute_sms above that; with fused shared experts it is additionally tightened by the combine geometry to

\max(\min(\texttt{compute\_sms},\,e+1),\,\lceil T\cdot\text{combine chunks}/8\rceil)

Panel of Figure 2 — Figure 2: Architecture of one GC-MoE layer replay. The child CUDA Graph holds three nodes: role1 (dispatch) and role2 (grouped GEMM) have no dependency edge and are concurrent roots, whereas role3 (finalize/combine) depends on both (solid arrows). Dashed arrows map each role to the green context that launches it—role1 to the 20-SM communication context, role2 and role3 to the 128-SM compute context—each with its own balanced workqueue so that the driver co-schedules them on disjoint SM sets. All cross-rank state lives in one symmetric-memory workspace whose boxes name the buffers exchanged over NVLink (double arrows): activations and scales, routing metadata, the role1 $\rightarrow$ role2 token ring, the role2 $\rightarrow$ role3 combine buffer, and the counters and parity signals. Every synchronization object is device-persistent and replay-aware, so a replay needs no host-side reset.

	Blackwell (SM103, B300)	Hopper (SM90, H20)
Compile target	sm_103f	sm_90a
MMA	tcgen05 mxf8f6f4.block_scale , 2-CTA	register-A WGMMA m64nNk32 , E4M3
Weight format	packed FP4 with UE8M0 K32 scales	MXFP4, window-repacked with per-expert factor
Activation format	E4M3 with packed UE8M0 K32 scales	E4M3FN with FP32 K128 scales
FP4 handling	native E2M1 operand (format 5)	in-register decode FP4 $\rightarrow$ E4M3
Scale application	UTCCP to TMEM; hardware block scale	per-K32 FMA outside MMA; $\times 64$ headroom

Table 2: SM103 (Blackwell) and SM90 (Hopper) implementation paths.

GPUs	8 $\times$ NVIDIA H20 (Hopper, sm_90a ); 8 $\times$ NVIDIA B300 (Blackwell, CC 10.3)
CUDA / assembler	CUDA 13.2, ptxas 13.2.78
PyTorch	2.11.0 (cu130)
Kernel DSL	NVIDIA CUTLASS DSL (CuTe) 4.8.1
Baseline	DeepGEMM 2.8.2+ fcad603 ( fp8_fp4_mega_moe on B300, fused shared experts when enabled; our SM90 reference port sm90_fp8_fp4_mega_moe on H20)
Proposed	GC-MoE 0.3.1 (commit 9bdb81f )
Execution	one process per GPU; PyTorch symmetric memory over NVLink; CUDA Graph replay

Table 3: Evaluation environment.

Model	$H$	$I$	$E$	top- $k$	shared experts	shared K group
GLM-5.2	6144	2048	256	8	1	128
DeepSeek-V4.1-Flash	5120	2304	384	6	1	32

Table 4: Model configurations. Routed recipe

(1,1,32)

for both models; shared-weight K group and shared recipe as shown.

Tokens/rank	Total tokens	Mega ( $\mu$ s)	GC-MoE ( $\mu$ s)	Speedup
1	8	120.061	115.161	1.0425
2	16	169.340	144.259	1.1739
4	32	256.018	202.589	1.2637
8	64	368.254	284.310	1.2953
16	128	442.722	341.689	1.2957
32	256	472.090	373.393	1.2643

Table 5: Decode latency and speedup, DeepSeek-V4.1-Flash on 8

\times

H20 (routed experts). Slowest-rank median per-invocation latency under CUDA Graph replay.

Figure 3: DeepSeek-V4.1-Flash routed-expert decode on 8 $\times$ H20: latency and speedup of GC-MoE versus Mega MoE.

Model	Tokens/rank	Total tokens	Mega ( $\mu$ s)	GC-MoE ( $\mu$ s)	Speedup
GLM-5.2	1	8	64.184	58.601	1.0953
	2	16	87.457	79.228	1.1039
	4	32	108.085	98.225	1.1004
	8	64	125.241	115.041	1.0887
	16	128	131.652	121.385	1.0846
	32	256	133.215	121.493	1.0965

Table 6: Decode latency and speedup, 8

\times

B300 (routed experts), both models. Slowest-rank median per-invocation latency. DSV4.1-Flash abbreviates DeepSeek-V4.1-Flash.

Figure 4: GLM-5.2 routed-expert decode on 8 $\times$ B300: latency and speedup.

Figure 5: DeepSeek-V4.1-Flash routed-expert decode on 8 $\times$ B300: latency and speedup.

Tokens/rank	gc-se ( $\mu$ s)	gc-unfused ( $\mu$ s)	unfused/se	mega-unfused ( $\mu$ s)	mega/se
1	145.171	156.866	1.0806	161.804	1.1146
2	179.286	188.566	1.0518	214.191	1.1947
4	236.353	248.009	1.0493	299.391	1.2667
8	302.667	328.262	1.0846	409.599	1.3533
16	363.640	384.421	1.0571	482.938	1.3281
32	428.722	422.456	0.9854	511.952	1.1941

Table 7: Decode latency, fused shared experts, DeepSeek-V4.1-Flash on 8

\times

H20. gc-se : GC-MoE with fusion; gc-unfused : GC-MoE with the shared expert run separately; mega-unfused : baseline. Ratios are against gc-se.

Figure 6: DeepSeek-V4.1-Flash with shared experts on 8 $\times$ H20: gc-se versus gc-unfused and mega-unfused.

Model	Tokens/rank	Total tokens	Mega-se ( $\mu$ s)	GC-MoE ( $\mu$ s)	Speedup
GLM-5.2	1	8	68.739	59.752	1.1504
	2	16	93.033	79.730	1.1669
	4	32	113.672	99.059	1.1475
	8	64	130.904	115.870	1.1297
	16	128	136.624	121.372	1.1257
	32	256	138.641	123.645	1.1213

Table 8: Decode latency and speedup with fused shared experts, 8

\times

B300. The baseline is Mega MoE in its fused shared-expert mode. Slowest-rank median per-invocation latency. DSV4.1-Flash abbreviates DeepSeek-V4.1-Flash.

Figure 7: GLM-5.2 with shared experts on 8 $\times$ B300: latency and speedup.

Figure 8: DeepSeek-V4.1-Flash with shared experts on 8 $\times$ B300: latency and speedup.

	TFLOPS		HBM (GB/s)		NVLink (GB/s)
Tokens/rank	Mega	GC	Mega	GC	Mega	GC
1	28.297	29.501	7 516.4	7 836.2	5.171	5.391
2	40.125	47.101	9 770.0	11 468.7	7.424	8.714
4	53.080	67.079	11 235.5	14 198.7	9.942	12.564
8	73.805	95.597	12 814.4	16 598.0	14.076	18.233
16	122.782	159.087	14 311.1	18 542.7	23.522	30.478

Table 9: Effective throughput and analytic bandwidth models, DeepSeek-V4.1-Flash on 8

\times

H20. TFLOPS counts matmul work only; HBM is the expert weight+scale working-set model; NVLink is the one-way assignment payload model. All values are aggregate over eight GPUs; GC denotes GC-MoE.

	GLM-5.2		DeepSeek-V4.1-Flash
Tokens/rank	Mega	GC	Mega	GC
1	75.280	82.453	48.101	57.208
2	110.496	121.974	80.869	90.891
4	178.817	196.765	129.099	142.488
8	308.642	336.007	194.199	212.781
16	587.226	636.895	336.480	366.062

Table 10: Effective TFLOPS (matmul work only, aggregate over eight GPUs), 8

\times

B300, routed experts. GC denotes GC-MoE.

Category	Cases	GC/BF16	Mega/BF16	GC/Mega	Max proc. diff
default	9/9	0.977039	0.977039	0.99999976	$1.80\times 10^{-5}$
boundary	23/23	0.977672	0.977672	0.99999976	$9.95\times 10^{-6}$
asymmetric	15/15	0.977837	0.977837	0.99999976	$1.80\times 10^{-5}$
GC: max absolute error 0.109375; max relative L2 0.00600836286321; repeat diff 0.
Mega: max absolute error 0.109375; max relative L2 0.00600836286321.

Table 11: Numerical accuracy, DeepSeek-V4.1-Flash on 8

\times

H20, routed experts. Thresholds: cosine

\geq 0.97

; processed diff and eager/graph repeat diff

\leq 0.002

. Cosine columns are per-category minima.

Appendix figures & tables1 asset

Supplementary material from the paper’s appendix.

Appendix

Constraint	Technical reason
$I\leq 4096$ , $I\bmod 128=0$	W1 output tiles $2I/128$ must fit the 64-bit l2_full_mask ; K128 tiles complete.
$H\bmod 256=0$	2-CTA clusters pair N128 tiles, requiring $H/128$ even.
$k+\mathbf{1}[\text{shared}]\leq 32$ , $k\geq 1$	One warp lane carries one combine slot; dispatch reads $32/k$ tokens per warp.
$E\leq 2048$ global, $\leq 512$ per rank, $E\bmod P=0$	Static per-expert signal region; lane-held per-expert counters.
$P\leq 72$ ranks	592-byte by-value peer descriptor with 72 offsets.
Routed recipe $(1,1,32)$ ; SwiGLU	MMA scale-factor path hard-codes K32 UE8M0; SwiGLU fused in the W1 epilogue.

Appendix

Table 12: Supported envelope and rationale (SM103 path unless noted).

Mixture-of-Experts Inference

Mixture-of-experts inference studies how to run sparsely activated expert models efficiently while preserving their quality.

See all 189 papers

Oct 7, 2026

Democratizing MoE inference on commodity GPUs with CoMoE

Deploying Mixture-of-Experts (MoE) models relies heavily on Expert Parallelism, which generates intense inter-GPU communication. Consequently, state-of-the-art inference systems require high-bandwidth, P2P interconnects (e.g., NVLink) in datacenter GPUs to handle massive token routing, making deployment prohibitively expensive. Consumer GPUs offer comparable compute power at significantly lower cost, promising to democratize MoE inference for individuals and enable privacy-preserving local deployments. However, their bandwidth-limited (only weak PCIe bus bandwidth) and host-mediated interconnects (no P2P support) introduce severe communication bottlenecks. We present CoMoE, a communication-efficient MoE inference system that resolves this mismatch through novel host-centric routing. Our key insight is that the unique communication topology provides the opportunity to elevate the host to an active routing hub, which can fundamentally reduce communication volume and eliminate global synchronization-induced stalls. Specifically, for token dispatch, we introduce host-backed token multicast to write shared tokens to the host exactly once, eliminating outbound transmission redundancy. For token combine, we propose a fine-grained, token-level aggregation mechanism using host staging buffers, which replaces rigid global synchronization and mitigates straggler effects. Evaluation on RTX 5090 GPUs shows that CoMoE improves inference throughput by up to 1.46x, approaching the performance of NVLink-capable A800 GPUs at only 23.4% of the hardware cost.

Ruwen Fan, Yuezhi Zu, Minhui Xie +6

Jimmy · Alibaba Cloud Computing · Tsinghua University

Communication-Efficient Mixture-of-Experts Mixture-of-Experts Inference

Ask PDF

Jul 21, 2026

Fine-grained Computation-Communication Overlap via Tile-level Signaling and Scheduling for Mixture-of-Experts

Mixture-of-Experts (MoE) architectures increase model capacity without proportionally increasing computation cost and have become a key building block for scaling large language models (LLMs) to trillion-parameter regimes. Efficient deployment of these MoE models relies on distributed execution across multiple GPUs, where each MoE layer involves two all-to-all communications: dispatching tokens to expert ranks and returning the expert outputs to their source ranks. Conventional MoE implementations launch this return all-to-all after expert compute completes, exposing communication latency on the critical path and reducing GPU utilization. We present a fine-grained approach that overlaps expert compute with the second all-to-all via tile-level signaling and scheduling. Our producer-consumer co-design combines: (1) a persistent per-rank computation kernel (producer) that covers all local experts on the rank to eliminate repeated kernel launch overhead and prioritizes remote-critical tiles, and (2) a persistent communication kernel (consumer) on a small dedicated partition of streaming multiprocessors (SMs) that issues segment-granular transfers as tiles become ready. Our co-design avoids intrusive changes to the underlying computation operators or communication primitives, making it practical for improving distributed MoE execution efficiency on multi-GPU systems. On a 4-A100 GPU platform, evaluated on three MoE models against four state-of-the-art MoE systems, our approach achieves up to 2.64x end-to-end speedup and 2.74x MoE-layer speedup. Compared with a conventional non-overlap baseline, our approach consistently improves both operator- and MoE-layer-level performance across varying GEMM shapes, router modes, and a broad range of producer/consumer SM partitions, while preserving correctness.

Minyu Cui, Anna Wingkvist, Morgan Ericsson

Linnaeus University · Växjö, Sweden

GPU Kernel Optimization Communication-Efficient Distributed Training Mixture-of-Experts Inference

Ask PDF

Mar 8, 2026

Scalable Training of Mixture-of-Experts Models with Megatron Core

Scaling Mixture-of-Experts (MoE) training introduces systems challenges absent in dense models. Because each token activates only a subset of experts, this sparsity allows total parameters to grow much faster than per-token computation, creating coupled constraints across memory, communication, and computation. Optimizing one dimension often shifts pressure to another, demanding co-design across the full system stack. We address these challenges for MoE training through integrated optimizations spanning memory (fine-grained recomputation, offloading, etc.), communication (optimized dispatchers, overlapping, etc.), and computation (Grouped GEMM, fusions, CUDA Graphs, etc.). The framework also provides Parallel Folding for flexible multi-dimensional parallelism, low-precision training support for FP8 and NVFP4, and efficient long-context training. On NVIDIA GB300 and GB200, it achieves 1,233/1,048 TFLOPS/GPU for DeepSeek-V3-685B and 974/919 TFLOPS/GPU for Qwen3-235B. As a performant, scalable, and production-ready open-source solution, it has been used across academia and industry for training MoE models ranging from billions to trillions of parameters on clusters scaling up to thousands of GPUs. This report explains how these techniques work, their trade-offs, and their interactions at the systems level, providing practical guidance for scaling MoE models with Megatron Core.

Zijie Yan, Hongxiao Bai, Xin Yao +42

NVIDIA

Low-Precision Training Mixture-of-Experts Models Communication-Efficient Distributed Training Memory-Efficient Training

Ask PDF

Abstract

Mixture-of-Experts Inference

Explore similar work

Democratizing MoE inference on commodity GPUs with CoMoE

Fine-grained Computation-Communication Overlap via Tile-level Signaling and Scheduling for Mixture-of-Experts

Scalable Training of Mixture-of-Experts Models with Megatron Core