arXiv CS open access

ISER-MoE: Dynamic In-Switch Replication of Hot Experts for Load Balancing in Expert-Parallel

Jintao Peng et al. · arxiv_cs

arXiv CS · Papers

cloud distributed-computing parallel-computing

Conceptio Open Archive

Analyze & explore related research in Conceptio

Connect this record directly into your research workflow or AI agent. Explore citations across open repositories, query full text, and export structured proof bundles.

Explore in Search →Connect via MCP / API ↗

ISER-MoE: Dynamic In-Switch Replication of Hot Experts for Load Balancing in Expert-Parallel Jintao Peng Ethan Zhang Kefeng Duan Bangyu Sheng Min Li Itamar Rosenfeld Rauch Jack Yang June Yang NVIDIA {jintaop, etzhang, kduan, bangyus, lmin, irauch, jackyang, juney}@nvidia.com

arXiv:2611.09008v1 [cs.DC] 6 Nov 2026

Abstract

routing, and the concurrent MegaFlux [38] does so inside the MegaMoE kernel that we also extend. Replication must move whole experts (37.16 MB each for DeepSeek-V4-Pro in NVFP4, NVIDIA’s 4-bit floating-point (FP4) format) within the layer, which raises three problems. Motivation 1 (Section 3.1): in a peak-rate model of FP4 prefill on GB300 (EP8), routed-expert computation hides only one or two DeepSeek-V4-Pro expert copies, whereas the shared expert, which needs no token traffic, leaves room for copies to all four helper slots, and on VR200 (Vera Rubin NVL72) for copies in 11 of 16 modeled scenarios instead of none. Motivation 2 (Section 3.2): a hot expert needs several replicas from one owner, so point-to-point fan-out makes the owner a hot spot. Motivation 3 (Section 3.3): baseline broadcasts pay latency or bandwidth overheads, from rings, from moving the payload twice in scatter + all-gather, or from the single injecting port of in-switch multicast. In-switch multicast, NVLink’s Scalable Hierarchical Aggregation and Reduction Protocol (NVLink SHARP, NVLS), removes the owner hot spot: the switch replicates each store into every member GPU [11, 17], so each byte of a hot expert leaves its owner once, whatever the replica count. Multicast groups, however, are fixed before use, and UltraEP set inswitch multicast aside because replica groups change across layers and micro-batches [36]. ISER-MoE (In-Switch Expert Replication for MoE) fixes the groups instead of the replicas (Figure 6). At each MoE layer invocation, its scheduler, HALO-Q (Hot-Aligned, Locality-Optimal expert replication with Quota repair), places each hot expert on one of a hierarchy of aligned groups created at initialization and reroutes part of its tokens to the replicas. A replication kernel then multicasts the expert from its owner or through SAMI (Scalable Asynchronous Multi-port In-switch broadcast), which spreads the injection over the members’ ports. Both run beside the shared expert, so the MoE kernel waits only for the plan. In trace replay, HALO-Q removed 98–99% of the excess load at EP4–24 (8,192 tokens/rank). On GB300 at EP8, context-only (prefill-only) throughput rose 17.4% on GSM8K and 7.2% on agentic prompts, more than with perfectly balanced routing (13.5% and 6.0%). In disaggregated serving on a GB300 NVL72 rack, it improved P90 request latency by 3.1–4.3% and output throughput by 1.4–1.9%. On pre-release Rubin GPUs at EP8, context-only throughput rose 6.1% on agentic prompts (BF16 combine).

Expert-parallel mixture-of-experts (MoE) layers run at the pace of the most loaded GPU, and a few hot experts carry most of the excess load. Replicating them at every MoE layer balances the load only if their weights arrive within the layer, and broadcasting from the owner is the bottleneck: point-to-point copies overload it, and baseline broadcasts pay latency or bandwidth overheads. We present ISER-MoE (In-Switch Expert Replication for MoE), built in TensorRT-LLM, which multicasts hot experts through the NVLink switch to fixed GPU groups, so each byte leaves its owner once. A GPU scheduler assigns hot experts to groups per layer, and a multi-port broadcast spreads the injection over the group’s ports, both overlapped with the shared expert. In trace replay, the scheduler removed 98–99% of the excess load; in prefill serving of DeepSeek-V4-Pro, throughput rose 17.4% on a math benchmark and 7.2% on agentic prompts.

Introduction

Agentic workloads send long, multi-turn prompts to large mixture-of-experts (MoE) models [1, 29], and under disaggregated serving [25, 44] their prefill holds most of the GPUs: published agentic configurations of DeepSeek-V4-Pro [7] on GB300 use 24 context (prefill) GPUs against 8 or 16 for generation [24]. Expert parallelism (EP) spreads each MoE layer’s routed experts over P GPUs (ranks; EP8 denotes P = 8), and fused MoE kernels such as MegaMoE [22, 42] synchronize all ranks, so each layer runs at the pace of its most loaded rank. On agentic coding traces, the rank imbalance ratio R of DeepSeek-V4-Pro prefill, the most loaded rank’s routed tokens over the mean, averaged 1.21–1.45 at EP4–8, and 1.69 at EP8 on GSM8K, the Grade School Math 8K benchmark (Section 2.2). Training and reinforcement-learning post-training face the same imbalance [2, 36, 45]; the method applies in principle to prefill, generation and training, and we implement and evaluate it for prefill. A few hot experts carry most of each layer’s excess (Section 2.3), so replicas in spare helper slots on underloaded ranks can absorb it. Static history-based plans [4, 5] remove standing skew but not its drift between batches, so recent systems replicate hot experts at every MoE layer invocation: UltraEP [36] and MoonEP [2] plan replicas on the GPU from the current

(b)

3 mean

IFBench

1.41

GSM8K

1.69

Agentic

1.45

SciCode

1.40

MMLU

1.42

Random

2.00 0

MoE layer

1.8 GSM8K 1.6

Agentic MMLU IFBench SciCode

1.4 1.2 4

EP degree

Figure 1: Rank imbalance ratio R (max/mean routed tokens per rank) of DeepSeek-V4-Pro prefill varies by workload, layer and EP degree (whole-trace loads, contiguous experts per rank). (a) EP8; right: 61-layer means. (b) 61-layer means. Workloads: GSM8K, MMLU (Massive Multitask Language Understanding), IFBench, SciCode, agentic coding traces and random tokens.

Background and Observation

2.1

Expert-Parallel MoE

MMLU

IFBench layer 60

Top-16: 42.45% Top-8: 31.47% Top-4: 22.24% Top-2: 14.68% Top-1: 8.81%

Top-16: 25.34% Top-8: 17.81% Top-4: 12.52% Top-2: 8.76% Top-1: 5.98%

Agentic

SciCode

100

Top-16: 30.22% Top-8: 21.78% Top-4: 14.90% Top-2: 9.44% Top-1: 5.73%

Mixed 20

Top-16: 32.11% Top-8: 21.42% Top-4: 13.42% Top-2: 7.94% Top-1: 4.45%

Top-16: 31.58% Top-8: 21.49% Top-4: 14.10% Top-2: 8.76% Top-1: 5.09%

Top-16: 29.40% Top-8: 19.92% Top-4: 13.02% Top-2: 8.07% Top-1: 4.69%

100 0

100

Experts sorted by load, descending (%)

Figure 2: Expert load is long-tailed (DeepSeek-V4-Pro, whole datasets): cumulative share of a layer’s assignments over its 384 experts sorted by load; thin: per MoE layer; thick: 61-layer average with the shares of the Top-1 to Top-16 experts; dashed: uniform load.

Our contributions are: • Multicast weight replication (Sections 4.1 and 4.2): to our knowledge, the first to replicate experts per invocation by in-switch multicast, at the cost of a full copy and a helper slot per member of a fixed aligned group. • SAMI (Section 4.3): the classic scatter + all-gather broadcast [32, 35] applied to weight replication, fused per slice, with the all-gather multicast from every member’s port. • A shared-expert copy window (Sections 3.1 and 4.4): a peak-rate model of FP4 prefill, after the DeepSeek-V4 report and the MoE inference system PROBE [7, 46], indicating that the shared expert gives copies the room routed experts lack, and a runtime that hides planning and replication in it. • Artifact: ISER-MoE in TensorRT-LLM as DynamicEPLB, a dynamic expert parallelism load balancer, in an open pull request under review [26] (Section 5).

GSM8K

100

Random

2.0

Cumulative share of assignments (%)

EP8 R

61-layer mean R

(a)

for each layer invocation from its current batch. In-switch multicast (NVLS) [11] replicates stores inside the NVLink switches: a store to a CUDA multicast object reaches the memory of every member [17]. Membership is fixed before use and only members map the object, so a writer is a member and receives its own replica; each object takes one of a limited number of switch multicast IDs (MLIDs; 1,024 per GB200 NVLink domain [16]). Streaming multiprocessors (SMs) issue copies through the Tensor Memory Accelerator (TMA) without the host; copy engines (CEs) need no SMs but are host-driven.

2.2

Rank Imbalance Varies by Workload and Layer

We recorded the prefill routing of DeepSeek-V4-Pro [7] (61 MoE layers, E = 384, k = 6), before load balancing, on five public workloads (GSM8K [3], MMLU [9], IFBench [27], SciCode [33] and agentic coding traces from SemiAnalysis [30]) and random tokens, and computed each workload’s wholetrace R per layer, rank r holding experts [rH, (r + 1)H). R depends on the workload: its 61-layer mean at EP8 was 1.40– 1.45 for four workloads, 1.69 for GSM8K and 2.00 for random tokens (Figure 1(a)). It depends on the layer (1.13–2.51 across the agentic traces’ layers at EP8), and any two workloads share at most one of their five most imbalanced layers. It grows with the EP degree, from 1.17–1.31 at EP4 to 1.40–1.69 at EP8 (Figure 1(b)). An invocation, however, sees only its batch: in a mixed trace interleaving the five workloads by request (4,340 requests, 18.2M prefill tokens, 91% agentic), R per iteration (forward pass) of 8,192 tokens/rank averaged 1.20, 1.44, 1.84 and 2.24 at EP4, EP8, EP16 and EP24, against 1.41 wholetrace at EP8. Most of a fixed mix’s skew is thus standing, and a static plan from load history, as in DeepSeek’s Expert Parallelism Load Balancer (EPLB) [5], removes most of it but not its drift between batches (Section 6.1). Imbalance thus costs time: on GB200 (EP8, 8,192 tokens/rank, trace-derived routing), MegaMoE without replicas took 513 + 1051 R µs

An MoE layer routes each token to k of its E experts [13, 31]. EP spreads the experts over P ranks, each holding H = E/P home experts; every invocation, in prefill, generation or training alike, dispatches tokens to their experts’ ranks and combines the outputs back. We target DEP, attention data parallelism (DP) with EP, in which every rank contributes its own tokens (tokens/rank, counted before top-k expansion). Many MoE models add a shared expert that processes every token [7]. Fused MoE kernels such as MegaMoE [22, 42] run dispatch, both expert general matrix multiplications (GEMMs) and combine in one persistent kernel that synchronizes all ranks, so every rank finishes with the most loaded one. Per layer, the load Lr of rank r counts the token–expert assignments (routes) it processes; the rank imbalance ratio R = maxr Lr /L̄ lies between 1 and P, R − 1 is the excess, and ranks above (below) L̄ are hot (cold). A replica copies a hot (heavily routed) expert into a helper slot, a per-rank buffer for one more expert, and takes routes from the expert’s owner, its home rank. Per-invocation planning [2, 36, 38] places replicas 2

(a) GB300, FP8 ×3: 12/16 → 16/16 above the line (16,667)

element). Each of the n expert copies a rank receives or sends moves W = 3He Ib bytes. Per link direction,

math/comm (FLOP/B)

30k 20k 10k 0

6 τkHe I math = . P−1 comm 2 P τkHe b + |{z} nW {z } |

(b) VR200, FP8 ×2: 0/16 → 11/16 above the line (23,333) 30k

dispatch + combine

20k

Communication hides behind computation at peak rates only if this ratio exceeds the device ratio F4 /Blink of dense FP4 FLOP/s to per-direction link bandwidth: 16,667 FLOP/B on GB300 (15 PFLOP/s, 900 GB/s [18]). On VR200 it is 23,333 FLOP/B (35 PFLOP/s and 1.5 TB/s per GPU and direction, from the rack datasheet [21]). Routed experts only. We model DeepSeek-V4-Pro (He = 7168, k = 6) and Kimi K3 [14], a latent MoE (top-16 of 896 experts, He = 3584); both have I = 3072 and are assumed NVFP4 (Kimi K3 exceeds our implementation’s 384expert limit and is only modeled). Without copies the ratio is 3I/b · P/(P − 1) = 18,725 FLOP/B at EP8, whatever τ, He and k; each copy lowers it, and of 16 scenarios (0, 1, 2 or 4 copies; 8,192 or 16,384 tokens/rank), 12 stay above GB300’s ratio (Figure 3(a), red). At 8,192 tokens/rank, DeepSeekV4-Pro’s routed computation (432.9 µs) exceeds its dispatch plus combine (385.4 µs) by only 1.15 copies of 41.3 µs; four copies would outlast the computation by 117.6 µs and delay the combine (Figure 3(c)); at 16,384 tokens/rank, two hide. Counting the shared expert. The shared expert needs neither the replica plan nor token communication, so it can run before dispatch, while the links are idle. Weighting its FP8 FLOPs by F4 /F8 = 3 on GB300 (assuming FP8 shared experts in both models), the ratio rises to 19,661–32,768 FLOP/B and all 16 scenarios allow overlap, four copies per rank included (Figure 3(a), green). On VR200, where FP8 runs at half the FP4 rate, the shared expert raises the count from 0 to 11 (Figure 3(b)), but DeepSeek-V4-Pro still hides at most one copy. On GB300, DeepSeek-V4-Pro’s shared expert lasts 216.5 µs at 8,192 tokens/rank, room for about five more copies at line rate. Copies should therefore start right after routing, beside the shared expert, not at the entry of MegaMoE (Figure 6(a)). Limits of the model. The model assumes peak rates, uniform routing and an NVFP4 combine; with the implementation’s default 16-bit brain floating-point (BF16) combine, token traffic alone outlasts DeepSeek-V4-Pro’s routed experts on GB300 (685 vs. 433 µs at 8,192 tokens/rank), so copies hide only beside the shared expert. It also counts helper tiles, which cannot hide their own copy (Section 4.4), as computation. At the 166–450 GB/s of the replication kernel on four GB200 SMs (Section 6.2), four copies take 330–898 µs and spill into the routed phase, which R1 allows (Section 3.4).

10k 0 n: 0 1 2 4 0 1 2 4 tokens/rank: 8,192 16,384 DeepSeek-V4-Pro

(c)

routed experts 4 copies

routed only

dispatch

+ shared expert

shared expert 4 copies

dispatch

1 2 4 0 1 2 4 8,192 16,384 Kimi K3 exposed combine routed experts combine

Figure 3: Motivation 1: with the shared expert, the copies overlap in all 16 modeled scenarios on GB300, against 12 with routed experts only (peak-rate model per MoE layer and rank; EP8, uniform routing, NVFP4, token traffic not deduplicated, as in MegaMoE). (a) GB300 [18]: math/comm ratio (FLOPs per communicated byte) with n copies per rank; red: routed experts only; green: with the 8-bit floating-point (FP8) shared expert at the FP4:FP8 rate (×3); hatched: below the device ratio (dashed). (b) VR200 [21], the same with ×2. (c) Example from (a) in time (DeepSeek-V4-Pro, 8,192 tokens/rank, four copies): beside the routed experts, the copies delay the combine; beside the shared expert, they hide.

over the 61 layers (Pearson correlation 0.938).

2.3

Observation: Long-Tail Distribution Enables Hot-Expert Replication

Expert load is long-tailed (Figure 2): over the mixed trace, the four busiest of 384 experts received 13.02% of a layer’s assignments on average (1.04% if uniform), and per iteration (EP8, 16,384 tokens/rank) the four busiest still received 13.6%. A hot rank’s excess thus sits on a few home experts: at EP8, the four busiest home experts of a layer’s hottest rank carried at least that rank’s excess in 59 of the 61 layers. Replicating them onto cold ranks can absorb most of the imbalance; HALO-Q therefore considers each rank’s min(S, H) busiest home experts (Section 4.1), with S = 4 helper slots per rank by default. Each replica, however, is a whole expert, copied within the invocation that uses it.

copies

Motivation

3.1 Motivation 1: More Weight-Copy Overhead Than Routed-Expert Computation Per-invocation replica copies [2, 36, 38] share the links with the layer’s token traffic; a peak-rate model of FP4 prefill shows how many of them computation can hide. Math/comm ratio. Per MoE layer and rank, τ tokens make τk token–expert pairs. Each costs 6He I FLOPs in its routed expert (He the expert’s input width) and, under uniform routing, leaves the rank with probability P−1 P , moving He b bytes in both dispatch and combine (b = 0.5625 B per NVFP4

3.2

Motivation 2: Source Broadcast Hot Spot

Hot experts often need several replicas each, and since they are a hot rank’s busiest home experts (Section 2.3), one owner exports several of them, while each receiver takes at most one per helper slot. If an owner replicates K experts, expert e to de ranks, point-to-point copies send ∑e de ≤ K(P − 1) payloads out of the owner; relays, as in UltraEP, which identified and 3

(a) Point to point: up to P−1 = 7 send-outs per hot expert

× 5×

hot spot

A A A A A B (c) Same greedy planner (simulation): max rank send-out

helper slot

A A A A A (b) In-switch multicast: 1 send-out per hot expert NVLink switch replicates

× 1× 0A

× 1× 5B

weights sent

helper slot

15 10 5 0

width falls to 250 GB/s at g = 16. B: software scatter + all-gather. The classic large-message broadcast [32] spreads injection over all g ports but moves the payload twice in series, TB = 2 g−1 g W /Blink + 2α, which caps its bandwidth near Blink /2. C: CE in-switch broadcast. A CE writer injects each byte once and the switch makes all g replicas, but from one ingress port replicated at a fixed rate: on four GB200 trays its total output stayed at 2.0–2.1 TB/s, so each receiver got 535, 266 and 127 GB/s at g = 4, 8 and 16. One injection per byte is thus necessary but not sufficient; injection must also be spread over many ports.

× 2×

point to point, no send limit 12 10

multicast

4 EP8

4 EP16

3.4

EP24

Design Requirements

The three motivations yield four requirements. R1 (Motivation 1): start planning and copying right after routing, hide them behind the shared expert and then the home-expert computation, and make only helper work wait for replica weights. R2 (Motivation 2): bound the experts each owner exports per layer, since even multicast sends each out of the owner once. R3 (Motivation 2): each byte of a hot expert leaves its owner once, whatever the replica count; as multicast membership is fixed before use and stores land at the same offset on all members (Section 2.1), this requires fixed groups and symmetric helper slots. R4 (Motivation 3): spread injection over many ports, by distinct owners or by group members, also when the owner lies outside the group. Sections 4.4, 4.1, 4.2 and 4.3 meet R1–R4 in order.

Figure 4: Replication makes the owner a broadcast hot spot. (a, b) EP8, idealized: hot experts A on GPU 0 and B on GPU 5 need five and two replicas, one per helper slot (below each GPU). (c) Simulated trace replay of a greedy planner (DeepSeek-V4-Pro, 8,192 tokens/rank, S=4), point to point or multicast: the busiest rank’s worst-case send-out per layer invocation.

modeled this fan-out bottleneck [36], relieve the owner at the cost of a store-and-forward hop but still inject ∑e de in total. In-switch multicast sends each byte out of the owner once, K payloads, with no extra hop (Figure 4(a, b)). We replayed the mixed trace (8,192 tokens/rank, S = 4; simulated assignment load) through a greedy point-to-point planner modeled on UltraEP and MegaFlux [36, 38]: like the placement stage of HALO-Q (Section 4.1), it moves the hottest rank’s busiest home experts to the coldest ranks, but copies each replica directly from its owner to any below-mean rank with a free slot. Unlimited, it balanced slightly better than HALO-Q (R = 1.001–1.009 against 1.006–1.020 at EP8–24), but its busiest owner sent 7/9/10 payloads per invocation at EP8/16/24 at the 99th percentile (Figure 4(c)), while no rank received more than four; at line rate, these 99th-percentile send-outs take 289–413 µs, longer than the 216.5 µs shared-expert window. Limiting each owner to four copies removed this tail but raised R to 1.016–1.100, whereas HALO-Q’s multicast plans kept the busiest owner at a median of 2/3/3 payloads, never above four.

ISER-MoE Design

Figure 6 shows one MoE layer under ISER-MoE: after routing, the HALO-Q scheduler decides which hot experts to replicate where and maps every route to an expert slot (Section 4.1), and a replication kernel multicasts the replicas into aligned rank groups (Sections 4.2 and 4.3), beside the shared expert (Section 4.4).

4.1

HALO-Q Scheduler

HALO-Q (Hot-Aligned, Locality-Optimal expert replication with Quota repair) is one GPU kernel per MoE layer invocation. It replicates hot experts in aligned rank groups (Section 4.2) with two objectives, in order (Figure 6(b)): • Objective 1, balance. MegaMoE finishes a layer when its most loaded rank does (Section 2.2), so HALO-Q brings every rank’s routed load down to the water level T . • Objective 2, locality. A route served off its source rank costs a token sent in dispatch and returned in combine, and the busiest sender or receiver bounds that traffic. Among balanced routings, HALO-Q therefore calculates as many tokens as possible locally. The five stages are (Algorithm 1): 1. Count exchange. Ranks exchange their E route counts through a symmetric NVLink buffer, with no host step, and

3.3 Motivation 3: Baseline Broadcasts Are Limited by Latency and Bandwidth Overhead Multicast lets each byte leave the owner once, but each of three ways to broadcast an expert to a group of g GPUs pays a latency or bandwidth overhead (Figure 5(a), A–C). We model the algorithm bandwidth of A and B from the per-direction link bandwidth Blink and a latency α per step, and measure C (Figure 5(b)). A: software ring. The NVIDIA Collective Communications Library (NCCL) 2.31.2 broadcasts only along a pipelined ring [10, 20]. Each link carries the payload once, but cbyte slices add a pipeline fill that grows with the group, TA = W /Blink + (g − 1)α + (g − 2)c/Blink , so the ring’s band-

A: Software ring broadcast

(b) Broadcast algorithm bandwidth

B: Software scatter + all-gather 0

C: CE in-switch broadcast NVLink switch

scatter

CE copy

all-gather Latency grows with the GPU count

Limited by 2× bandwidth overhead

0 1 2 3 Limited by one ingress port’s multicast rate

algorithm BW (GB/s)

(a) Baseline broadcasts (g=4)

1000

Blink (900 GB/s)

750 500

B: scatter + all-gather (model)

250

A: ring (model) C: one port (measured)

group size g (GPUs)

Figure 5: Baseline broadcasts are limited by latency or bandwidth. (a) One expert’s weights, in four slices, from GPU 0 to g=4 GPUs. (b) Algorithm bandwidth (BW; payload/time) of broadcasting to g GPUs: A and B modeled for one expert (37.16 MB) with Blink = 900 GB/s, α = 5 µs and 2-MiB ring slices; C measured on four GB200 trays (160 MiB).

Algorithm 1 HALO-Q on EP rank s, once per MoE layer invocation; all ranks compute the same plan from the same all-gathered counts. Spare quota: quota beyond own routes.

partner’s own routes), quota for its own routes of an expert both host, and hands the partner as much of its own spare quota; run in both orders, this cancels opposite flows as far as the loads allow, while every load and total, hence balance and copy count, stays. 5. Route rewrite (local first). Without communication, each rank first fills its own quota; the surplus routes of all ranks, in rank order, then fill the other hosts’ spare quotas in rank order (a north-west-corner assignment), following UltraEP’s local-first cumulative quotas [36] and MoonEP’s prefix sums [2]. A route runs in its host’s helper slot, or in the home slot if the host is an owner outside the group. Replicas are byte copies and gate weights are untouched, so replication changes where a route runs, not what it computes.

Require: routes Rs ; P ranks with H home and S helper slots; aligned groups Ensure: broadcasts b = (eb , Gb , hb ); a slot per route (home unless rerouted) 1: C ← A LL G ATHER(counts of Rs ); loads L; T ← ⌈∑r L[r]/P⌉ ▷ stage 1 2: while a rank above T can still export do ▷ stage 2: Objective 1 3: o ← the busiest such rank; e ← its heaviest unexported home expert 4: G ← the aligned group, at the finest level able to absorb e (else coarser), with a free slot on every member and the most room below T ; if none has room, freeze o 5: otherwise b ← (e, G) with quotas qb : move up to L[o] − T of e’s routes from o to G \ {o}, coldest first, each up to T 6: hb ← the first index free on all of Gb , coarse levels first ▷ stage 3 7: for all rank pairs (a, c), twice, in both orders do ▷ stage 4: Objective 2 on experts both host, a trades spare quota with c in equal amounts, as 8: much as both allow: it takes what its own routes need and gives what they do not 9: for all broadcasts b do ▷ stage 5 10: s keeps up to qb [s] of its routes to eb ; the surplus routes of all ranks fill the hosts’ spare quotas in rank order 11: a route runs in helper slot hb of its host (the home slot if the host is the owner outside Gb )

4.2

Multi-Level Symmetric Multicast

Replicating a moderately hot expert on all P ranks wastes helper slots, yet arbitrary groups cannot be prepared: multicast membership is fixed before use, MLIDs are limited (Section 2.1), and the ranks needing a replica change with every layer and batch. ISER-MoE therefore fixes the groups, not the replicas: level ℓ splits the P ranks into 2ℓ aligned groups of gℓ = P/2ℓ consecutive ranks, for ℓ = 0 and every finer level with 2ℓ | P and gℓ ≥ 4 (e.g., 24, 12 and 6 ranks at EP24), and HALO-Q picks a level per expert (Figure 6(d)). The four-rank floor is a software choice, not a multicast limit: imbalance falls as each rank holds more experts (Section 2.2), and smaller groups are unevaluated. All groups’ multicast objects are created at initialization. Each rank holds its H = E/P home experts in local slots [0, H) and S helper slots after them. A broadcast installs an expert’s weights and block scales (its bundle) in one helper slot on every member of an aligned group. Symmetric helper slots. A multicast store lands at the same offset on every member, so a broadcast needs one helper index free on all members. Alignment guarantees one: if every rank belongs to the groups of at most S broadcasts, coloring broadcasts from coarse to fine, each with the smallest index free on all its members, never fails. Aligned groups are nested or disjoint, so every earlier, equally coarse or coarser group that meets a group G contains it, and all members of G see the same at most S − 1 used indices. HALO-Q ensures the precondition by admitting a broadcast only if every member has a free slot. Without alignment, S+1 pairwise-overlapping

run the same deterministic planner, so all obtain the same plan without broadcasting it, as in UltraEP [36]. 2. Placement (Objective 1). With T = ⌈∑r Lr /P⌉, HALO-Q repeatedly takes the busiest rank above T that can still export and its heaviest home expert not yet exported. An owner exports at most S experts per invocation, chosen among its min(S, H) heaviest, which bounds its egress (Section 3.2); a rank receives at most S replicas. The expert goes to the aligned group, not necessarily containing the owner, whose members all have a free helper slot, at the finest level that can absorb the expert at the water level (coarser only if none there has room), with the most room below T ; up to the owner’s excess over T of its routes move there, coldest member first, each up to T . An owner for which no group has room, or which has no candidate left, is frozen. 3. Slot coloring. Each placed expert gets one helper index free on all members of its group (Section 4.2). 4. Quota repair (Objective 2). Placement fixes loads and each replicated expert’s total quota, not its split among hosts, so rank a may send tokens of one expert to rank c while c sends tokens of another to a. In two rounds of a circle-method tournament over disjoint rank pairs, each rank takes, from its partner’s spare quota (beyond the 5

(a) Per-layer schedule on one rank (DEP) MAIN stream

event / ready flag COPY-stream kernel MAIN-stream kernel READY gate: all ranks at γ wait MegaMoE (fused kernel) quantize expert computation: home tiles helper tiles

shared expert

router

in-switch weight replication (TMA) per slot: owner-direct or SAMI

HALO-Q scheduler

HALO-Q next layer

READY(γ) notification

Objective 1, balance: every rank’s routed load → water level T 1600 1 other experts 1300 2 ✓ all ranks at T 600 hot experts A, B 300 1000 1000 1000 1000 T B 700 B replica routes A A B A above T : moved 400 1 2 before

reuse fence

token combine

(b) Zoom into HALO-Q: a two-objective greedy scheduler

consumer done(γ)

metadata router + token dispatch

plan ready expert IDs COPY stream (highest priority)

time

NVLink switch

s3 s2

s1 ingress egress

1 2 3 after placement

ingress egress

Objective 2, locality: min token send-out / receive-in: calculate tokens locally after placement: tokens cross B tokens

quota repair 3 AB 1 AB trade equal A tokens quota: A only quota: B only A and B quota

after repair: no tokens cross loads unchanged

TMA copy

✓

GPU 0

1 AB calculates own B tokens

TMA copy

GPU 1 GPU 2 GPU 3 Pipelined step 1 pull slice from owner GPU 0

3 AB calculates own A tokens

TMA copy

GPU 1 GPU 2 GPU 3 Pipelined step 2 multicast slice with all ingress ports

GPU 0

every member’s port injects; W /g extra inbound traffic per non-owner member

(d) Multi-level symmetric multicast (EP16, S=4) P P/2 P/4 h3 h2 h1 h0

level 0: one group of all 16 ranks level 1: {0–7} level 2: {0–3} D A 0

D A 1

D A 2

{8–15} {4–7}

D A 3

C A 4

C A 5

C A 6

{8–11} C A 7

B A 8

B A 9

B A 10

{12–15} B A 11

B A 12

B A 13

B A 14

B A 15

A, very hot: all 16 ranks (P). B: P/2. C, D: P/4, 4 slots each, not 16. Multicast objects (MLIDs, address ranges) per group, not per replica: 1 + 2 + 4 = 7 groups. HALO-Q picks the level per expert.

Figure 6: ISER-MoE in one MoE layer (schematic). (a) One rank (DEP): after routing, the highest-priority COPY stream runs HALO-Q and the replication kernel beside the shared expert; MegaMoE waits only for the plan, and only its helper tiles wait, at the READY gate, for generation γ (a per-invocation counter) from every rank. (b) HALO-Q (ranks 0–3). Objective 1, balance: the busiest rank’s hottest expert gets a replica (steps 1, 2), and the coldest ranks take its routes above the water level T , the mean rank load. Objective 2, locality: after placement, ranks 1 and 3 would exchange tokens of the multicast experts B and A; quota repair trades equal parts of the two quotas, so both calculate their own tokens locally, at unchanged loads. (c) Replication with SAMI (g = 4, owner GPU 0) in two pipelined steps: each member pulls its slice from GPU 0 (s1) and stores it through its own ingress port (s2), and the switch multicasts it to all members (s3). (d) Multi-level symmetric multicast: symmetric groups of P, P/2 and P/4 ranks, each with multicast objects; HALO-Q picks the level per expert. Rows h0–h3: helper slots.

groups would need S+1 indices.

For an expert of W bytes, the owner still sends W and each member injects W /g, so the output fed by one port, which limits single-port multicast (Section 3.3), falls from gW to W and the issue work is shared; the price is W /g of extra inbound traffic per non-owner member. Since a multicast writer must be a member, SAMI also serves owners outside the group, which only serve the pulls. A slot is replicated owner-direct if its owner belongs to the group G and either G is pure (every broadcast that meets G targets exactly G with its owner inside) or this is the owner’s widest broadcast into a group containing it; otherwise SAMI replicates it. This default rule avoids SAMI’s extra inbound traffic where the owner can serve its group, relying on distinct owners to spread the injection; whether it pays off with many owners is untested (Section 6.2).

4.3 TMA Weight Replication and SAMI Multi-Port Broadcast The replication kernel copies with TMA from SMs, avoiding the host round trip of CEs: warps stage 8-KiB slices in shared memory and store them to the group’s multicast address (multimem.cp.async.bulk), so the switch writes each slice into every member’s helper slot. Owner-direct replication has the owner, a group member, multicast the whole bundle through its own port. With SAMI (Figure 6(c)), the g members share the injection in two pipelined steps, with no barrier or reassembly: member i (s1) pulls each slice j ≡ i (mod g) from the owner and (s2) stores it through its own port to the multicast address, and (s3) the switch delivers it to every member. SAMI is thus a scatter + all-gather broadcast [32, 35] whose all-gather is in-switch multicast, as in NVLS AllGather [20].

4.4

Runtime Overlap

(a) Objective 1, balance: per-layer imbalance ratio

1.75

no balancing

HALO-Q, S = 4

EP4

1.20 1.002 (1.008)

1.50

2.5

EP8

1.44 1.006 (1.021)

EP24

2.24 1.020 (1.063)

1.5

1.00 4

HALO-Q, S = 3

2.0

1.25

R (max/mean)

ISER-MoE issues the copy right after routing, where Section 3.1 found room beside the shared expert (Figure 6(a)): • The COPY stream, the model’s highest-priority stream, runs HALO-Q and the replication kernel while the MAIN stream computes the shared expert and quantizes the routed input. • MAIN then waits on the device for the plan and launches MegaMoE on the rewritten routes; dispatch never waits for the copy. • MegaMoE claims tiles in slot order, so dispatch overlaps the home tiles, and combine the later home and the helper tiles. • Only helper tiles wait: after its writes drain, the replication kernel multicasts the invocation’s READY generation to every rank, even with nothing to copy, and each cooperative thread array (CTA) waits before its first helper tile until every rank’s generation is current. To first order, the copy is exposed only if it outlasts the shared expert and the home tiles; a copy issued inside MegaMoE, as described for MegaFlux [38], would miss the shared-expert window.

1.0 EP16

6 5 4 3 2 1

1.84 1.010 (1.036)

3 2 1 0

20 40 60 0 20 40 layers (sorted by R without balancing); labels: 61-layer mean

routed tokens (×103 )

(b) Objective 2, locality: max token send-out / receive-in (S = 4) 100 75 50

max send-out max receive-in −8%

−19%

−6%

−31%

−3%

−45%

−2%

−55%

25 0

EP4

EP8

EP16

EP24

Figure 7: The two objectives of HALO-Q in simulated trace replay (DeepSeek-V4-Pro, 8,192 tokens/rank). (a) Objective 1, balance: per-layer R without balancing, layers sorted by R (dotted: R = 1); labels: 61-layer means without balancing and with HALO-Q, S = 4 (S = 3). (b) Objective 2, locality (S = 4): the most routed tokens any rank sends out or receives in during dispatch, mean per iteration and layer; labels: change from no balancing.

Implementation

ISER-MoE is implemented in TensorRT-LLM as DynamicEPLB, in an open pull request under review [26]; we describe its commit 99dac8f1. A model enables this per-layer dynamic mode through moe_config.load_balancer, with num_slots = P(H + S) and an SM budget, auxiliary_sms. Kernels, streams and SM budgets. We extend TensorRTLLM’s NVFP4 MegaMoE kernel [22], written in the CuTe domain-specific language, with helper slots and the READY gate. Its per-expert weight and scale planes append, after the H home rows, one of two helper banks shared by all layers (layer i uses bank i mod 2); as in UltraEP and MoonEP [2, 36], sharing cuts helper storage to 297 MB per rank at S = 4 (computed), from 9.07 GB for 61 per-layer slot sets. Before each HALO-Q launch, COPY waits until MegaMoE has consumed the previous replicas; since every plan starts with the all-rank count exchange, no copy overwrites a bank another rank still reads. All MoE layers share the COPY stream, whose two kernels run back to back: HALO-Q on auxiliary_sms CTAs (default 8) and the replication kernel on 8. The budget does not partition the GPU: the routed-input quantizer uses the remaining SMs, and MegaMoE the full grid. The shared expert is the standard gated multilayer perceptron (MLP) with no SM budget, so its overlap with COPY rests on stream priority. An SM-budgeted fused shared expert exists only for one other GPU generation. Because persistent MegaMoE CTAs spin at the READY gate, the copy must obtain SMs while MegaMoE runs; this rests on scheduling (early enqueue, highest priority, 8 independent CTAs), not on a hardware guarantee, and a starved copy would hit the gate’s 120-s timeout and fail rather than hang. Stale or failed plans, malformed copy plans, gate errors and timeouts abort the job, with no fallback to home experts.

Scope. The mode requires the MegaMoE NVFP4 backend, DEP, an even EP degree up to 32, E ≤ 384 divisible by P, CUDA 13.1 or later and one multicast-capable NVLink domain. CUDA graphs are unsupported, because the runtime binds one MAIN submission stream, orders COPY with events and keeps per-submission host state. No test launches HALO-Q or the replication kernel, and the generation phase is untested.

Evaluation

Methodology. All experiments use DeepSeek-V4-Pro’s MoE layer with NVFP4 routed experts (Section 2.1). A GB200 compute tray holds four GPUs, so EP8 spans two trays. OFF denotes a layer without helper slots and ON the same layer with S helper slots per rank. Unless stated otherwise, kernel times are the slowest rank’s median over 20 iterations, reported as the time reduction (OFF−ON)/OFF or the speedup OFF/ON.

6.1

E1: Simulated Scheduler Rank Imbalance

E1 replays the mixed trace of Section 2.2 at 8,192 tokens/rank through the placement, quota repair and local-first route rewrite of HALO-Q, planning each (iteration, layer) independently; it counts assignments, not time. A CPU reference of the HALO-Q kernel and an earlier GB200 build of it matched the simulator bit for bit; the kernel at commit 99dac8f1 has not been checked.

one broadcast at a time, slowest rank; best of each

SAMI (multi-port)

2.03× 3.76×

200 single port

CE prototype

600

481

486

400 258

200

SAMI owner-direct

173

600

TMA copy, 4 SMs, EP8

400 200

8 GPUs 12 GPUs 16 GPUs

copy SMs per GPU

142

1.34×

36.3

30 1.14×

14.7

1.03×

2.0

0 1

group size g (GPUs)

(d) time increase (%)

1.13×

8 GPUs, 283.5 MiB, one copy at a time

400

tuned ncclBroadcast (nccl-tests)

600

single-port (owner-direct) multicast

(b) algorithm BW (GB/s)

algorithm BW (GB/s)

SAMI

(a)

284

message size (MiB)

0.6

2.9

16K

0.3

32K

tokens/rank

Figure 8: Reference data (E2): broadcasts from one owner on GB200. (a) CE prototype: algorithm bandwidth of SAMI (multi-port) and single-port multicast vs. group size g (g = 4: two trays, 144 MiB; g = 8, 16: four trays, 160 MiB); labels: SAMI’s speedup. (b) A development build of the replication kernel (four SMs by default), standalone, slowest rank: algorithm bandwidth vs. copy SMs per GPU, eight 37.16 MB bundles (283.5 MiB, (c)’s largest size) from one owner to 8 GPUs, the owner included. (c) One broadcast at a time, until the slowest rank holds it: the fastest SAMI variant (mostly with the owner outside the group) vs. the fastest tuned ncclBroadcast setting (nccl-tests, NCCL 2.31.2: best of protocol, buffer size, channel count, buffer registration and CUDA-graph launch); line style: GPU count. (d) Time increase from the shared expert’s start to the end of MegaMoE when the build of (b) copies four experts (4 × 36 MiB) during the shared expert (EP8); labels: SAMI’s speedup. (a)

3,000

skewed routing (power law 0.8)

3,000 2,000

balanced routing OFF ON

1.21×

L20 (R = 2.36): 2,843 → 1,626 µs, 1.75×

1.22×

kernel time (µs)

4,000

(b)

1.52× 1.24×

1,000

2,500 mean 1,999 µs

2,000

1.23× 1,624 µs

1,500 0

EP4 8K

EP8 8K

EP4 16K

EP8 16K

OFF

balanced routing

time saved

MoE layer

Figure 9: MegaMoE on GB200 (DeepSeek-V4-Pro layer, NVFP4, FP4 combine, four helper slots; slowest rank, median of 20 iterations; replicas and routes from an earlier all-rank policy, installed before timing). (a) Power-law routing (exponent 0.8); 8K and 16K are tokens/rank; labels: speedup. (b) 61 layers, routes sampled from the trace (EP8, 8,192 tokens/rank); shading: time saved; margin: means; L20: layer 20. Dashes: balanced routing (OFF).

Objective 1, balance. With S = 4, HALO-Q lowered the 61-layer mean of R from 1.20/1.44/1.84/2.24 to 1.002/1.006/1.010/1.020 at EP4/8/16/24, removing 98.4– 99.0% of the excess load (S = 3: 94.9–95.9%), and kept every layer within 13% of balance (Figure 7(a)). An earlier policy that replicates four experts on every rank, which also produced the routes of E3, left 1.038/1.193/1.348 at EP8/16/24 with 28/60/92 remote copies per iteration and layer, against HALO-Q’s 20.4/41.4/76.2. Objective 2, locality. The most routed tokens that any rank sent out fell by 2–8%, and the most that any rank received in by 19–55% (Figure 7(b)). Balance does most of the latter, as hot owners no longer receive the excess; quota repair alone cut both maxima by a further 0.5–2.3%. Static plan. A static per-layer plan with the same four extra slots per rank, built by our re-implementation of the EPLB rule [5] from past or whole-trace loads under an optimistic even token split, left R = 1.06–1.07 at EP8, against 1.006 with per-invocation planning. By the GB200 fit of Section 2.2

(513 +1051 R µs), extrapolated below its fitted range and without helper-tile overhead, a static plan would capture 85–88% of the kernel-time gain of per-invocation planning over no balancing. Alignment and SAMI. A broadcast charges a helper slot on every group member; the chosen group excluded the owner in 53% of broadcasts at EP8 and 76% at EP16–24. Only SAMI serves such groups (Section 2.1); restricted to groups that contain the owner, R rose to 1.038/1.201/1.296 at EP8/16/24. Per-rank point-to-point placement. The point-to-point planner of Motivation 2 (Section 3.2) avoids that charge: with at most four copies per owner and four replicas per rank, it made 2.5–3.3× fewer remote copies than HALO-Q (EP8–24) but balanced worse.

6.2

E2: SAMI Broadcast (Reference Data)

E2 compares SAMI with the single-port and ring broadcasts of Section 3.3, using a CE prototype, which issues the steps of Section 4.3 as host CE copies of one contiguous share 8

per member, and two other builds of the replication kernel in standalone harnesses. One owner broadcasts bundles to g GPUs, itself included except in (c)’s TMA build. Against single-port multicast. In our placements, a single injecting port lost per-receiver bandwidth as g grew, whereas SAMI stayed near the owner’s egress limit (Figure 8(a)): its speedup grew from 1.13× at g = 4 to 2.03× at g = 8 and 3.76× at g = 16. At g = 4, HALO-Q’s group size at EP8 and EP16 in replay, one port’s 535 GB/s per receiver is already close to that limit. SM-issued copy. The implementation’s TMA copies add a per-SM issue limit. Figures 8(b, d) use a development build of the replication kernel (four SMs by default). Standalone with eight bundles (283.5 MiB; Figure 8(b)), owner-direct multicast was issue-bound, rising from 46 GB/s on one SM to 258 GB/s on eight, whereas SAMI, with every member issuing a share, reached 481 GB/s from four SMs on (with four bundles, 166 and 450 GB/s on four SMs). The default per-slot rule (Section 4.3) nonetheless picks owner-direct here: 258 GB/s on eight SMs. Against ncclBroadcast. Figure 8(c) completes each broadcast on every rank before the next starts and compares the fastest tuned ncclBroadcast setting of nccl-tests with the fastest SAMI variant, mostly another TMA build of the replication kernel with the owner outside the multicast group (Section 4.3). We tuned NCCL 2.31.2 at every point over the protocol, the buffer size, the channel count, buffer registration and CUDAgraph launch; the LL128 protocol with CUDA-graph launch won all 12 points and cut the default ring’s time (Simple protocol, 512-KiB slices) by 12–70%. SAMI was 1.07/1.21/1.34× faster for one expert (35.4 MiB) at 8/12/16 GPUs and 1.22– 1.29× for eight: it took 101–107 µs per expert at every g, against 108–143 µs for the tuned ring, which still grows with g (the default ring by 20 µs per added GPU). For eight experts on 8 GPUs, this variant reached 591 GB/s against 486 GB/s in (b), where the owner, a group member, also receives every replica, which slows the pulls it serves; (b)’s SAMI matched tuned ncclBroadcast there (1.00×). Overlap with the shared expert. Figure 8(d) runs the fourexpert copy (EP8; balanced synthetic routing, dummy weights) during a 7168→2048→7168 FP8 shared expert, narrower than the model’s. With SAMI, the copy took 310–330 µs and lengthened the span by 2.0/0.6/0.3% at 8K/16K/32K tokens/rank; with owner-direct multicast, it took 0.98–1.20 ms and lengthened the span by 36.3/14.7/2.9%, mostly where the copy outlasted the shared expert (256/480/936 µs; Section 3.1). Because every rank copies a share, SAMI slowed the shared expert by 10/4.8/2.0% and MegaMoE by at most 0.5%. With the replication kernel of 99dac8f1 on eight SMs, SAMI slowed the 8K shared expert by 49%: in the profile, the copy ran beside the shared expert’s GEMMs for 99.9% of its time rather than in launch gaps, and the GEMMs, with no SMs set aside for the copy, slowed instead; the SM-free CE prototype slowed it by 2.9/1.8/0.9%. End to end on prerelease Rubin GPUs (EP8), a fused shared expert that leaves the copy’s four SMs free ran 156–161 µs per layer beside the copy, against 153–158 µs without copies, and the copy

spent 71–80% of its time beside it and 2.5–5.6% in idle gaps (medians per rank).

6.3

E3: MegaMoE Kernel Time

We time one call of a later MegaMoE revision than 99dac8f1’s for a DeepSeek-V4-Pro layer on GB200, OFF and ON with four helper slots per rank, under two routings: synthetic power-law routing at four shapes (EP4 and EP8; 8,192 or 16,384 tokens/rank), each with a balanced control, and per-layer routes from the trace. All runs used FP4 combine instead of the default BF16 combine. ON’s replicas, routes (from E1’s all-rank policy) and READY flags were installed before timing, so the times exclude HALO-Q, the copy and any READY wait. Each cell or layer ran as OFF, ON, OFF (OFF: the mean of its two runs); outputs matched bitwise. Power-law routing. With power-law popularity (exponent 0.8), the replicas reduced the skewed cells’ R from 1.27–1.69 to 1.00–1.03 and kernel time by 17.2–34.1% (1.21–1.52×), to within 0.3–3.5% of balanced routing (Figure 9(a)); on balanced routing, time changed by −0.7% to +0.4%.1 Trace routing. For each of the 61 MoE layers, we sampled EP8 routes at 8,192 tokens/rank from the layer’s expert popularity in the mixed trace, calibrated to its EP8 imbalance within 3% (one seed). Every layer improved, from 1.3% (layer 45, R = 1.12) to 42.8% (layer 20, R = 2.36), and the mean time fell from 1,999.3 to 1,624.3 µs (1.23×), 2.1% above balanced routing (Figure 9(b)): OFF time grew linearly with R (Section 2.2), whereas ON time was uncorrelated with it (Pearson 0.01). Both routings are stationary, so E3 measures the kernel’s gain from replicas, not the value of per-invocation planning (Section 6.1). ON’s slowest layers carry the largest imbalance left by all-rank replication, which HALO-Q reduces in replay only (worst EP8 layer 1.066 vs. 1.198). Costs outside the kernel. In a sequential GB200 run of an earlier build, with ranks released together, HALO-Q took 19.6– 29.2 µs and the replication kernel (8 SMs, including READY) 128.7–340.6 µs; serialized, the copy would cost 31–63% of the kernel time saved on skewed routing (derived across runs; indicative), so it must overlap other work. Planning time grows with the EP degree: on one CTA, placement and quota repair took 3.8, 6.1, 8.7 and 12.2 µs at EP4, 8, 16 and 32 (medians on the trace’s routing counts, three helper slots, four at EP4; pre-release Rubin GPU), as the greedy placement makes about 19 sequential exports per invocation at EP32 against 5 at EP8. In an overlapped run of an earlier build beside a 300 µs singleCTA delay kernel standing in for the shared expert, with ranks arriving 10 µs apart, ON took 15.6–29.2% less time than OFF on the skewed cells and 0.2–1.5% more on the balanced ones, planning, copy and READY waits included. The stand-in moves no data, so this run misses interference with the real shared expert. 1We report medians because single outlier iterations can dominate the mean of 20, the runs’ declared statistic; with the mean, the skewed cells took 17.4–34.0% less time and the balanced cells changed by −5.0% to −0.3%, the extremes due to single slow OFF iterations.

ON (ISER-MoE) FB (force expert balance)

17.4

1200 Output throughput (tok/s/GPU)

E2E throughput gain over OFF (%)

13.5 9.5

KV reuse 94.4–94.6%

10.8

7.2

6.0

KV reuse 94.9–95.0%

3.8 1.6

DEP8

DEP6

DEP8

GSM8K

DEP6 AgentPerf

Gain over OFF (%)

4.2

3.4

3.9

2.9

16K DEP8

3.9

16K DEP6

Figure 11: Context-only benchmark on pre-release Rubin GPUs (DeepSeek-V4-Pro NVFP4, attention DP; 8K and 16K: at most 8,192 and 16,384 tokens/rank per iteration): gain of ON over OFF in computed-token throughput on AgentPerf prompts (whiskers: the reported ±). Every configuration gained.

6.4

C3: 48 GPUs P90 +3.1%, tput +1.6%

400

C4: 40 GPUs P90 +4.3%, tput +1.9%

35 40 45 50 P90 request end-to-end latency (s)

ON (ISER-MoE) OFF 55

bine, throughput rose 6.1% and 9.1% at EP8 (8,192 and 16,384 tokens/rank) and 4.2% and 3.9% at EP6 (±0.6–1.1 points; Figure 11), and time to first token (TTFT) improved 6.6–12.1% except at EP6 with 8,192 tokens/rank (−1.2 ± 2.4%); with FP4 combine, throughput rose 2.6–3.9%. A profile of an earlier commit (one rank, EP8, BF16 combine, S = 3) shows the routed MegaMoE kernel 29.5% and all MoE work 15.0% shorter per layer. The shared expert, beside the copy and with SMs held back for it, slowed by 127 µs (from 182.6 µs); the profile does not separate the two causes.

2.6

600

Figure 12: E5 on GB300 NVL72: disaggregated serving of DeepSeek-V4-Pro (NVFP4) on agentic coding traces, output throughput per GPU against P90 request end-to-end latency, OFF and ON (ISER-MoE), configurations C1–C4 joined in order of concurrency. Labels: improvement of ON, OFF/ON−1 for latency and ON/OFF−1 for throughput (tput); each arm’s best run.

6.1 5

800

C2: 40 GPUs P90 +4.1%, tput +1.4%

BF16 combine, 6 runs FP4 combine, 2 runs

9.1

1000

200

Figure 10: E4 on GB300: end-to-end context-only throughput gain over OFF of ON (ISER-MoE) and of FB (force expert balance: OFF with routing forced to be perfectly balanced); DeepSeek-V4-Pro NVFP4, tile-wise FP4 combine. AgentPerf labels: key-value (KV) cache reuse rate across the three arms.

C1: 56 GPUs P90 +4.0%, tput +1.6%

6.5

E4: End-to-End Prefill Serving

E5: End-to-End Disaggregated Serving

We served DeepSeek-V4-Pro (NVFP4) in TensorRT-LLM with prefill and decode disaggregated on a GB300 NVL72 rack, in four configurations: C1 to C4 ran five, three, two and one DEP8 prefill instances with one DEP16 (C1, C2) or DEP32 (C3, C4) decode instance, on 56, 40, 48 and 40 GPUs at 2,626, 1,152, 736 and 388 concurrent requests. Decoding used speculative decoding with three draft tokens and acceptance fixed at 2.01. Clients replayed agentic coding traces [30] with aiperf for a 1,200 s window after warm-up; ON used three helper slots per rank, and KV cache reuse was 94.6–95.2%, equal in both arms of every configuration. In C3 and C4, ON also kept replicas resident: a replication whose whole group already holds the expert in that helper slot moves no data; C4 also skipped replications that would move fewer than 256 routes. C1–C3 split the prefill instances between the arms, which shared the decode instance, in each of four, four and two jobs on the same nodes, alternating positions; C4 ran each arm in its own job, as two ON/OFF pairs on different nodes. Because serving performance varies between runs, each arm repeated one fixed configuration and reports its best run. ON shortened P90 request latency by 4.0, 4.1, 3.1 and 4.3% (OFF/ON−1) and raised output throughput per GPU by 1.6, 1.4, 1.6 and 1.9% in C1–C4 (Figure 12), so every configuration moved toward lower latency and higher throughput.

We served DeepSeek-V4-Pro (NVFP4) context-only in TensorRT-LLM with attention DP at EP8 or EP6 (DEP8, DEP6), OFF and with ISER-MoE (ON); throughput counts computed (non-reused) tokens per GPU and second. On GB300, a third arm, force expert balance (FB), ran OFF with TensorRT-LLM’s perfect router, which forces synthetic, perfectly balanced routing; its outputs are not meaningful. All GB300 arms used tile-wise FP4 combine, in which the second GEMM’s epilogue quantizes each output tile to NVFP4. On GSM8K [3], ON raised end-to-end throughput by 17.4% at DEP8 and 9.5% at DEP6, against FB’s 13.5% and 10.8% (Figure 10). On AgentPerf [1], with 94.4–95.0% key-value (KV) cache reuse in every arm, ON gained 7.2% and 1.6%, against FB’s 6.0% and 3.8%. These gains come from load balance; multicast keeps the busiest owner at four or fewer send-outs per layer, where point-to-point copies under the same greedy placement need up to 10–14 at EP8–24 (Section 3.2); the implementation has no unicast copy path, so we did not compare the two end to end. On pre-release Rubin GPUs, OFF and ON ran with at most 8,192 or 16,384 tokens/rank per iteration. On AgentPerf requests (95.1–95.9% KV cache reuse), throughput was measured in six runs per arm with BF16 combine, the implementation’s default, and in two with FP4 combine. With BF16 com-

6.6

Limitations

by predicting routing [46]; Megatron-Core and Kimi K3 overlap the shared expert with token communication [12, 37], but, to our knowledge, no prior system that replicates experts per invocation uses it to hide planning and replication. Reported gains. Reported gains are not directly comparable: MegaFlux reports 1.45× at kernel level, including planning and transfers that our 1.21–1.52× excludes, and 1.13–1.26× end to end in vLLM prefill of DeepSeek-V4-Pro; UltraEP reports 1.29× over SGLang with EPLB at EP40–64. In-network collectives. NVLS accelerates collectives over fixed groups [11, 20] and its AllGather and two-shot all-reduce already inject from every member’s port [15, 20]; SAMI applies this to weight replication, with an owner that may lie outside the group. MoE systems multicast plan metadata [2], tokens [19], activations [12] and bucket collectives [43] in the switch, and DySHARP proposes switch support for token multicast [41]; none multicasts expert weights.

On GB200 we have not measured a layer with the real shared expert or end-to-end throughput, and no hardware result compares with a static plan (Section 6.1). Each member of an aligned group spends a helper slot and receives the whole expert, and owner-direct replication was issue-bound on 8 SMs in a development build of the replication kernel (E2). Without a shared expert, only the home tiles can hide the copy, which leaves room for at most about one DeepSeek-V4-Pro-sized copy (GB300 peak-rate model, EP8, 8,192 tokens/rank) and none with a BF16 combine (Section 3.1). The scope is prefill with attention DP in one NVLink domain, in eager mode; the implementation accepts up to EP32 (Section 5), but hardware runs cover EP4–8 and replay EP4–24, and no controlled hardware point lies below 8,192 tokens/rank. At EP32, planning takes 12 µs, twice its time at EP8 (Section 6.3). Replicas live for one invocation, although in replay 74–78% of the all-rank policy’s expert selections recur in the next iteration; placement cannot fix ranks that reach the layer at different times.

Conclusion

Per-invocation replicas of hot experts can remove most of the rank imbalance in expert-parallel MoE layers, but only if their weights arrive within the layer. ISER-MoE multicasts them through the NVLink switch to aligned GPU groups created once, so each byte leaves its owner once; SAMI spreads the injection over the group’s ports, and planning and replication overlap the shared expert. In trace replay, HALO-Q removed 98–99% of the excess load, and on GB300 at EP8, contextonly throughput rose 17.4% on GSM8K and 7.2% on agentic prompts. On pre-release Rubin GPUs at EP8, context-only throughput rose 6.1% on agentic prompts. A hardware comparison with static plans, generation and training, attention-side imbalance and replication across NVLink domains remain future work.

Related Work

Static replication. Deployed systems replicate hot experts from load history: DeepSeek-V3 periodically adjusts redundant prefill experts [4], EPLB plans replicas from estimated loads [5], and serving engines move weights between rebalances [23, 34]. Such plans can worsen imbalance as load shifts [36]; the Linear-Programming-Based Load Balancer (LPLB) [6] re-splits each batch’s tokens over static replicas without moving weights, so it corrects drift only for replicated experts. METRO [39] routes over static replicas to minimize activated experts in memory-bound decode, not computebound prefill. Static plans and per-invocation planning are complementary (Section 6.1). Per-invocation replication. MegaFlux [38], concurrent with our work, replicates hot experts at every invocation inside the MegaMoE kernel that we extend and pipelines the transfer with per-projection readiness, hiding an estimated 56–76% of it; it describes neither multicast nor relays, so by our reading an owner sends one transfer per replica. UltraEP [36] rebalances every layer and micro-batch from the exact load, identifies hot-expert fan-out as the weight-distribution bottleneck, relays experts with many replicas in two stages, and sets in-switch multicast aside because replica groups change; we keep that premise and fix the groups instead, as FasterMoE broadcasts to one fixed group [8]. MoonEP [2], described in the Kimi K3 report [12], balances training load exactly with E/P prefetch slots per rank, each owner pushing weights with one TMA store per destination; ECHO (Elastic Cloning for Hot Experts) and the concurrent DeepEP V2.5 communication library also replicate experts dynamically [37, 43]. All of them copy replica weights point to point; ISER-MoE multicasts them and accepts a small residual imbalance. Hiding the copy. FEPLB and TAOT (Topology-Aware Optimal Transport) hide replica copies behind home-expert computation [28, 40], and PROBE behind earlier computation

References [1] Artificial Analysis. AA-AgentPerf methodology. Benchmark methodology web page, 2026. URL https://artificialanalysis.ai/me thodology/agentperf. [2] Y. Chen, C. Li, Y. Wang, and M. Wei. MoonEP: A perfectly balanced expert parallelism library via dynamic redundant experts. GitHub repository, 2026. URL https://github.com/MoonshotAI/MoonEP. [3] K. Cobbe, V. Kosaraju, M. Bavarian, M. Chen, H. Jun, L. Kaiser, M. Plappert, J. Tworek, J. Hilton, R. Nakano, C. Hesse, and J. Schulman. Training verifiers to solve math word problems, 2021. arXiv:2110.14168. [4] DeepSeek-AI. DeepSeek-V3 technical report, 2024. arXiv:2412.19437. [5] DeepSeek-AI. EPLB: Expert parallelism load balancer. GitHub repository, 2025. URL https://github.com/deepseek-ai/EPLB. [6] DeepSeek-AI. LPLB: Linear-programming-based load balancer. GitHub repository, 2025. URL https://github.com/deepseek-ai/LPLB. [7] DeepSeek-AI. DeepSeek-V4: Towards highly efficient million-token context intelligence, 2026. arXiv:2606.19348. [8] J. He, J. Zhai, T. Antunes, H. Wang, F. Luo, S. Shi, and Q. Li. FasterMoE: Modeling and optimizing training of large-scale dynamic pretrained models. In Proceedings of the 27th ACM SIGPLAN Symposium on Principles and Practice of Parallel Programming (PPoPP ’22), pages 120–134. ACM, 2022. doi: 10.1145/3503221.3508418. [9] D. Hendrycks, C. Burns, S. Basart, A. Zou, M. Mazeika, D. Song, and J. Steinhardt. Measuring massive multitask language understanding. In International Conference on Learning Representations (ICLR), 2021.

[10] Z. Hu, S. Shen, T. Bonato, S. Jeaugey, C. Alexander, E. Spada, J. Dinan, J. Hammond, and T. Hoefler. Demystifying NCCL: An indepth analysis of GPU communication protocols and algorithms, 2025. arXiv:2507.04786. [11] A. Ishii and R. Wells. The NVLink-network switch: NVIDIA’s switch chip for high communication-bandwidth superpods. In 2022 IEEE Hot Chips 34 Symposium (HCS), pages 1–23, 2022. doi: 10.1109/HCS559 58.2022.9895480. [12] Kimi Team. Kimi K3: Open frontier intelligence, 2026. arXiv:2607.24653. [13] D. Lepikhin, H. Lee, Y. Xu, D. Chen, O. Firat, Y. Huang, M. Krikun, N. Shazeer, and Z. Chen. GShard: Scaling giant models with conditional computation and automatic sharding. In International Conference on Learning Representations (ICLR), 2021. [14] Moonshot AI. Kimi K3. Hugging Face model card, https://huggin gface.co/moonshotai/Kimi-K3, 2026. [15] NVIDIA. 3x faster AllReduce with NVSwitch and TensorRT-LLM MultiShot. NVIDIA Technical Blog, 2024. URL https://develope r.nvidia.com/blog/3x-faster-allreduce-with-nvswitchand-tensorrt-llm-multishot/. [16] NVIDIA. NVIDIA GB200 NVL partition user’s guide. DU-12143001_0.3, Section 4.1.4.3, 2025. URL https://docs.nvidia.com/ multi-node-nvlink-systems/partition-guide-v1-2.pdf. [17] NVIDIA. CUDA driver API: Multicast object management. CUDA Driver API Reference Manual, Section 6.29, 2026. URL https://do cs.nvidia.com/cuda/cuda-driver-api/group__CUDA__MULT ICAST.html. CUDA 13.4. [18] NVIDIA. NVIDIA GB300 NVL72. Product page, 2026. URL ht tps://www.nvidia.com/en-us/data-center/gb300-nvl72/. Accessed 2026-10-12. [19] NVIDIA. Megatron-Core MoE inference token dispatcher (NVLSAllGatherVDispatcher). Source code, 2026. URL https://github.com /NVIDIA/Megatron-LM. megatron/core/transformer/moe/to ken_dispatcher_inference.py. [20] NVIDIA. NCCL: Optimized primitives for collective multi-GPU communication. GitHub repository, 2026. URL https://github.com/N VIDIA/nccl. [21] NVIDIA. NVIDIA DGX Vera Rubin NVL72 datasheet. Datasheet 5697900, preliminary, Sept. 2026. URL https://resources.nvid ia.com/en-us-dgx-systems/dgx-vera-rubin-nvl72-datash eet. Accessed 2026-10-14. [22] NVIDIA. CuTe DSL MegaMoE kernels in TensorRT-LLM. Source code, 2026. URL https://github.com/NVIDIA/TensorRT-LLM /tree/b5011b14415d45c2f0cfeb31cf50e2ab68377867/tenso rrt_llm/_torch/cute_dsl_kernels/cutedsl_megamoe. [23] NVIDIA TensorRT LLM Team. Scaling expert parallelism in TensorRT LLM (Part 1: Design and implementation of large-scale EP). TensorRTLLM technical blog, 2025. URL https://github.com/NVIDIA/Te nsorRT-LLM/blob/main/docs/source/blogs/tech_blog/blo g04_Scaling_Expert_Parallelism_in_TensorRT-LLM.md. [24] NVIDIA TensorRT LLM Team. DeepSeek-V4 on NVIDIA Blackwell: Model-specific and agentic-workload optimizations in TensorRT LLM. TensorRT-LLM technical blog, 2026. URL https://github.com/N VIDIA/TensorRT-LLM/blob/main/docs/source/blogs/tech_ blog/blog26_DeepSeek_V4_on_NVIDIA_Blackwell_Model_Sp ecific_and_Agentic_Workload_Optimizations_in_TensorR T-LLM.md. [25] P. Patel, E. Choukse, C. Zhang, A. Shah, Í. Goiri, S. Maleki, and R. Bianchini. Splitwise: Efficient generative LLM inference using phase splitting. In 2024 ACM/IEEE 51st Annual International Symposium on Computer Architecture (ISCA), pages 118–132, 2024. doi: 10.1109/IS CA59077.2024.00019. [26] J. Peng. Add DynamicEPLB kernels and runtime integration. TensorRTLLM pull request #19646, 2026. URL https://github.com/NVIDI A/TensorRT-LLM/pull/19646. Described at commit 99dac8f1744b (2026-10-12); squashed on 2026-10-13 as 81348989ba76, with byteidentical DynamicEPLB kernels. [27] V. Pyatkin, S. Malik, V. Graf, H. Ivison, S. Huang, P. Dasigi, N. Lambert, and H. Hajishirzi. Generalizing verifiable instruction following. In Advances in Neural Information Processing Systems (NeurIPS), Datasets and Benchmarks Track, 2025. [28] S. Qi, H. Liu, and S. Zhao. FEPLB: Exploiting copy engines

for nearly free MoE load balancing in distributed training, 2026. arXiv:2604.19654. [29] SemiAnalysis. AgentX methodology and datasets (InferenceX). Benchmark methodology web page, 2026. URL https://inferencex.s emianalysis.com/datasets. AgentX v1.0. [30] SemiAnalysis. CC Traces – Weka, With Subagents, v7 only. Hugging Face dataset, 2026. URL https://huggingface.co/datasets/se mianalysisai/cc-traces-weka-062126. [31] N. Shazeer, A. Mirhoseini, K. Maziarz, A. Davis, Q. V. Le, G. E. Hinton, and J. Dean. Outrageously large neural networks: The sparsely-gated mixture-of-experts layer. In International Conference on Learning Representations (ICLR), 2017. [32] R. Thakur, R. Rabenseifner, and W. Gropp. Optimization of collective communication operations in MPICH. International Journal of High Performance Computing Applications, 19(1):49–66, 2005. doi: 10.117 7/1094342005051521. [33] M. Tian, L. Gao, S. D. Zhang, X. Chen, C. Fan, et al. SciCode: A research coding benchmark curated by scientists. In Advances in Neural Information Processing Systems (NeurIPS), Datasets and Benchmarks Track, 2024. arXiv:2407.13168. [34] vLLM Project. vLLM v0.23.0 release notes: EPLB for DeepSeek-V4 Mega-MoE. GitHub release, 2026. URL https://github.com/vll m-project/vllm/releases/tag/v0.23.0. [35] G. Wang, S. Venkataraman, A. Phanishayee, J. Thelin, N. Devanur, and I. Stoica. Blink: Fast and generic collectives for distributed ML. In Proceedings of Machine Learning and Systems (MLSys), 2020. arXiv:1910.04940. [36] X. Wei, C. Jin, T. Dai, Y. Zhong, S. Yu, C. Yang, B. Wu, Z. Zhang, J. Mai, Q. Zhu, Z. Li, Y. Liu, and G. Luo. UltraEP: Unleash MoE training and inference on rack-scale nodes with near-optimal load balancing, 2026. arXiv:2606.04101. [37] Z. Yan, H. Bai, X. Yao, D. Liu, T. Liu, H. Liu, et al. Scalable training of mixture-of-experts models with Megatron Core, 2026. arXiv:2603.07685. [38] J. Yao, S. K. S. Hari, V. Balaji, S. Damani, I. Jang, P. Viswanath, and C. Kozyrakis. MegaFlux: Skew-resilient MoE megakernels via pipelined expert replication, 2026. arXiv:2610.00671. [39] Y. Yu, H. Ma, K. Agarwal, N. Oswald, Q. Huang, H. Linsenmaier, C. Mei, R. Zhao, R. Borkar, B. D. Rouhani, D. Nellans, R. Krashinsky, and A. Khandelwal. Efficient MoE serving in the memory-bound regime: Balance activated experts, not tokens, 2025. URL https://arxiv.or g/abs/2512.09277. arXiv:2512.09277. [40] L. Zhang, H. Zhang, S. Gu, K. Mo, S. Han, S. Li, Y. Wang, and D. Shen. TAOT: Topology-aware optimal transport for dynamic expert replica placement in MoE training, 2026. arXiv:2608.03676. [41] Q. Zhang, C. Zhang, Z. Zhou, H. Wang, Z. Zhou, Z. Tu, G. Sun, Z. Xie, Y. Diao, Z. Ji, J. Leng, G. He, and M. Guo. Accelerating MoE with dynamic in-switch computing on Multi-GPUs. In Proceedings of the 53rd Annual International Symposium on Computer Architecture (ISCA), 2026. arXiv:2605.05607. [42] C. Zhao, Z. Xu, L. Zhao, J. Li, C. Xu, A. Xu, S. Liu, K. Zhou, and K. Yu. DeepGEMM: Clean and efficient BLAS kernel library on GPU. GitHub repository, 2025. URL https://github.com/deepseekai/DeepGEMM. Mega MoE fused kernel added 2026-04-16. [43] C. Zhao, S. Zhou, L. Zhang, C. Deng, Z. Xu, Y. Liu, K. Yu, J. Li, and L. Zhao. DeepEP: An efficient expert-parallel communication library. GitHub repository, 2025. URL https://github.com/deepseekai/DeepEP. V2.5, commit def8651 (2026-09-29). [44] Y. Zhong, S. Liu, J. Chen, J. Hu, Y. Zhu, X. Liu, X. Jin, and H. Zhang. DistServe: Disaggregating prefill and decoding for goodput-optimized large language model serving. In 18th USENIX Symposium on Operating Systems Design and Implementation (OSDI 24), pages 193–210. USENIX Association, 2024. [45] Y. Zhou, H. Li, S. Lin, Y. Zhao, T. Zhao, X. Miao, J. Jiang, F. Fu, and B. Cui. Harnessing routing foresight for micro-step-level MoE load balancing in RL post-training, 2026. arXiv:2606.11867. [46] Q. Zhu, X. Ye, Y. Liu, H. Ouyang, and C. Song. PROBE: Co-balancing computation and communication in MoE inference via real-time predictive prefetching, 2026. arXiv:2602.00509.

ISER-MoE: Dynamic In-Switch Replication of Hot Experts for Load Balancing in Expert-Parallel

Analyze & explore related research in Conceptio

Related documents