ExpertPlex: A High-Goodput Disaggregated Serving System for MoE LLMs with Adaptive Persistent Kernels
arXiv:2607.18002v2 [cs.DC] 21 Jul 2026
Bingyang Wu1 Chao Jin1 Zili Zhang1 Xinming Wei1 Yinmin Zhong1 Ruidong Zhu1 Chengxu Yang2 Xin Jin1 Yuliang Liu2 1 Peking University 2 Independent Researcher Abstract latency-sensitive decode phase [13, 23, 38, 39, 58]. For better performance, PDD also needs to allocate resources according to the P:D ratio that matches the resource demands of two phases. However, existing instance-level PDD provisions a complete model replica for each phase, so its allocation granularity grows with the model. A reported DeepSeek-V3 deployment uses 32 GPUs for prefill and 320 GPUs for decode in one unit [13] to achieve this ratio. Small clusters cannot realize this ratio, leading to resource waste. Large clusters can achieve this ratio, but suffer from a larger failure blast radius because a rank failure in hierarchical communication can stall the entire unit [15, 44]. They can also scale only at the granularity of the large deployment unit when traffic shifts. Duplicated expert weights also displace KV cache. PDD therefore obtains isolation at the cost of memory efficiency, elasticity, and fault containment. Prefill-decode colocation solutions instead use Green Context [36] to reserve SMs for each phase to achieve weight deduplication and computation multiplexing [10, 22, 41], but they cannot follow the dynamic demands of MoE LLMs. When Green Context is used, resources are fixed for a kernel, and reconfiguration is limited at the prefill layer level because it requires CPU intervention. However, MoE and attention modules in a layer and communication and computation operations in MoE modules have different demands [43, 60]. Dynamic expert computation also creates layer-to-layer changes in expert activation on each GPU [20, 21, 47]. Coarse reconfiguration therefore creates head-of-line blocking and resource bubbles (§2). Partitioning a GPU also reduces the local resources available to each phase, increasing its required degree of parallelism with more communication while leaving cross-phase network interference unmanaged. To address these problems, our key insight is a hybrid disaggregation-colocation architecture that shares experts across phases but disaggregates their attention modules. MoE weights contribute over 95% of model parameters. Sharing eliminates their cross-phase memory duplication and multiplexes dynamic per-rank expert loads, letting computation from either phase fill bubbles in the other’s attention–expert pipeline. Attention holds under 5% of parameters, so disaggregation enables independent per-phase allocation in singleGPU units without duplicating massive expert weights. Because attention is more compute-intensive [43, 60], giving each phase whole attention GPUs rather than intra-GPU partitions preserves local compute capacity for each phase in a GPU. Partitioning attention GPUs would require wider
LLMs scale Mixture-of-Experts (MoE) parameters for superior intelligence, but massive weights and dynamic computation impede efficient serving. Existing instance-level prefill-decode disaggregation isolates the phases on separate full-model replicas. As MoE weights grow, each instance may span tens to hundreds of GPUs, making resource allocation increasingly coarse. Configured prefill-to-decode ratios thus often mismatch demand, overprovisioning one phase while overloading the other. Prefill-decode colocation avoids this duplication, but existing Green Context solutions partition each GPU by phase and fix phase resources during a kernel. They cannot track resource changes across operations or layerwise variation in routed expert load, causing head-of-line blocking or idle reserved resources. Partitioning every GPU also leaves each phase with fewer local resources, forces wider parallelism and more communication, and lets prefill and decode traffic interfere on the shared network. We present ExpertPlex, which shares massive MoE experts across phases while disaggregating lightweight attention modules. Expert sharing eliminates over 95% of duplicate model weights and multiplexes dynamically sparse computation, while attention disaggregation reduces attention communication cost. ExpertPlex further uses (1) adaptive persistent kernels to schedule dynamic expert computation at tile granularity for efficient, isolated execution; (2) attentioninitiated MoE communication to avoid network interference and enable cross-phase communication-computation overlap; and (3) a tile-to-cluster model to optimize these mechanisms for maximum goodput. Experiments serving MiniMaxM2.7 and GLM-5.1-FP8 show that ExpertPlex improves goodput by up to 2.01× over instance-level prefill-decode disaggregation and 1.66× over prefill-decode colocation.
1
Introduction
Scaling laws continue to reward larger model capacity [8, 26]. Frontier LLMs therefore increasingly use sparse Mixtureof-Experts (MoE) modules that dynamically activate a few experts to add parameters without proportionally increasing computation [13, 17, 27, 33, 40, 51]. It is still challenging to serve them efficiently, because instances must still store massive expert weights, while token-dependent routing makes expert computation and communication vary across tokens. Modern LLM serving systems typically use prefill-decode disaggregation (PDD) to serve MoE LLMs to avoid interference between the compute-intensive prefill phase and the 1
CTA Cluster
per-phase parallelism for equivalent resources, increase communication, and introduce cross-phase attention network interference. This architecture lets ExpertPlex match phase demands with fewer GPUs and less traffic, improving utilization, scaling granularity, and fault containment. However, sharing the experts requires finer GPU control than existing mechanisms provide. CUDA stream priorities [35] and API interception [49] cannot preempt longrunning prefill kernels that may significantly block latencysensitive decode kernels, while Green Context [36], MPS [1], and MIG [37] cannot reallocate resources at fine granularity between phases, leading to resource bubbles. Shared experts also couple the phases through communication. The traffic flows from two phases to the same MoE module can interfere with each other, and conventional two-sided MoE communication requires MoE-side coordination that can cause deadlock when different ranks are used by different phases. To address these challenges, we propose ExpertPlex, a disaggregated serving system for MoE LLMs with adaptive persistent kernels. ExpertPlex runs an Adaptive Persistent Kernel (APK) on each MoE GPU. APK schedules MoE computation at tile boundaries, providing bounded preemption for urgent decode work and reallocating idle CTA clusters to prefill without CPU intervention or kernel relaunch. The bound is independent of sequence length, while persistent execution preserves CUDA Graph compatibility and GPU utilization. ExpertPlex further uses attention-initiated onesided MoE communication to use mostly disjoint network paths to transfer activations between attention and final MoE-side buffers without APK coordination. It therefore removes MoE-side polling, avoids cross-phase deadlock and interference, and overlaps communication from one phase with computation from the other. Because these mechanisms introduce coupled choices in placement, parallelism, computation-communication overlap, and tile scheduling, ExpertPlex proposes a cross-stack placement optimizer that models them jointly to maximize goodput. Evaluations serving MiniMax-M2.7 and GLM-5.1-FP8 on real-world workloads show that ExpertPlex improves goodput by up to 2.01× over instance-level prefill-decode disaggregation, and 1.66× over Green Context-based prefill-decode colocation. This paper makes the following contributions. • We introduce a hybrid architecture that eliminates crossphase MoE weight duplication, multiplexes dynamically sparse expert computation, and disaggregates attention to reduce the per-phase degree of parallelism with less communication while preserving fine-grained isolation. • We design APKs for tile-level preemption and reallocation and attention-initiated one-sided communication to reduce network interference and enable cross-phase overlap between communication and computation. • We jointly optimize them across the stack and evaluate ExpertPlex on two frontier MoE LLMs, demonstrating goodput improvements over state-of-the-art baselines.
CTA i+1
CTA i
Math warp TMA warp
TMA warp
Tensor Cores
CUDA Cores
Math warp
Math warp
TMA Units
TMA warp
TMA warp
Tensor Cores
CUDA Cores
L1 Cache/SMEM
Math warp
TMA Units
L1 Cache/SMEM
L2 Cache Global Memory
Figure 1. GPU execution model.
2
Background and Motivation
2.1
MoE LLM Inference
An LLM stacks Transformer layers, each with an attention module and a feed-forward network (FFN) module. Attention mixes information across tokens and materializes keyvalue tensors as the KV cache, which later iterations reuse. The FFN instead transforms each token independently. Frontier MoE LLMs replace conventional FFNs with mixture-ofexperts (MoE) modules and scale their expert weights aggressively [13, 27, 33, 40, 51]. These weights now dominate the model footprint, such as 95% in DeepSeek-V4-Pro [3], 96% in GLM-5.1-FP8 [6], and 98% in MiniMax-M2.7 [5]. An MoE module partitions the FFN into many experts. Shared experts process all tokens, while a router dynamically selects a small top-𝑘 subset of routed experts for each token at each layer. The router dispatches the token activation to the activated experts and combines their outputs before the next layer. This sparse activation increases model capacity without proportionally increasing computation per token. Inference has two phases with different objectives. Prefill processes all input tokens in parallel, builds their KV cache, and produces the first output token, which makes it throughput-oriented. Decode then generates one token per iteration from the preceding token based on the prefix KV cache. It is latency-sensitive because decode latency is uservisible and orders of magnitude shorter than prefill latency. 2.2
GPU Execution Model
MoE modules are primarily grouped general matrix multiplications (GEMMs) on GPUs, so their performance depends on the GPU execution model. As Figure 1 shows, a GPU contains many streaming multiprocessors (SMs). Each SM integrates Tensor Cores, CUDA cores, Tensor Memory Accelerator (TMA) units, registers, an L1 cache, and software-managed shared memory. All SMs share the L2 cache and global memory, which is also visible to peer GPUs and NICs. Recent GPUs allow a thread-block cluster to coordinate across SMs through distributed shared memory (DSMEM) and to multicast a global-memory tile through TMA. 2
Larger DoP! (1.5x for prefill, 2x for decode) Prefill: 2/3
CPU Intervention Different Resource Demand (Use Different Green Contexts)
GPU 0
Decode: 1/3 Prefill: 2/3
GPU 1
Decode: 1/3
Prefill Attn
Prefill MoE Attn
Network Interference!
Prefill: 2/3
Prefill MoE
Prefill Attn
GPU 2
Decode: 1/3
Launch decode
Blocking!
MoE 5 Active Experts
MoE 2 Attn Active Experts
Bubble!
Dynamic Computation
Figure 2. Limitations of prefill-decode colocation solutions. GPU work is launched as kernels through CUDA streams. Kernels in one stream execute in order, while kernels in different streams may overlap when resources are available. Within a kernel, threads form warps, warps form cooperative thread arrays (CTAs), and CTAs may form clusters. A CTA runs on one SM until completion, while CTAs in a cluster can cooperate across SMs. High-performance GEMM kernels specialize warps for data movement and Tensor Core computation, then pipeline them to hide memory latency [15, 56]. 2.3
multicast over its faster scale-up fabric, avoiding redundant scale-out transfers when several activated experts reside on that node. Decode moves far less data per iteration, so direct transfers to the relevant ranks reduce latency. PDD also specializes expert computation [56]. Prefill packs many expert inputs into a contiguous layout to reduce memory. Decode has too few tokens to amortize this packing, so a masked layout avoids copies. Limitations. The isolation, however, creates a rigid and increasingly coarse deployment granularity. Each prefill or decode instance must hold a complete model replica, so growing MoE weights directly increase its minimum GPU count, and duplicated weights displace KV-cache capacity. PDD must then provision these indivisible instances in a prefill-todecode resource ratio that matches the phases’ asymmetric demands. The smallest combination that realizes this ratio defines the deployment unit, which becomes much larger than either phase’s instance alone. One reported DeepSeekV3 unit combines 32 prefill and 320 decode GPUs to realize its target ratio [13]. Another uses 176 GPUs as a single unit for a different load [14]. Kimi-K2 has likewise been deployed on 128 H200 GPUs [34]. This compounding effect makes the target ratio impossible to realize in a small cluster, while deviating from it makes one phase overloaded and leaves the other phase overprovisioned. For a large cluster with sufficient GPUs, each scaling step changes capacity by hundreds of GPUs and cannot match moderate traffic shifts efficiently. Larger units also widen the failure blast radius. Unlike direct transfers, hierarchical communication couples many ranks into one group, so a single rank failure can stall the entire unit and is difficult to recover from [15, 44]. PDD therefore obtains isolation at the cost of memory efficiency, elasticity, and fault containment.
Large-scale MoE LLM Serving
MoE weights often exceed one GPU’s memory, so serving systems use expert parallelism (EP) to shard experts across GPUs. Because routing decisions vary by token, dispatch sends activations to the ranks hosting the selected experts and combine returns their outputs. Both operations are commonly implemented as all-to-all communication [15]. Serving systems can hide part of this cost with two-batch overlap (TBO), which overlaps one microbatch’s communication with another’s computation, or single-batch overlap (SBO), which overlaps communication with shared-expert computation in the same microbatch. Attention admits a different parallelization because its weights are much smaller. Data parallelism (DP) replicates attention to favor throughput and KV-cache capacity with little communication. Tensor parallelism (TP) shards attention to reduce per-request latency, but adds communication and may reduce effective KV-cache capacity when entries are replicated. Therefore, the resource demand and best configurations of attention and MoE often differ [43, 60]. 2.4
Prefill-Decode Disaggregation
Prefill-decode disaggregation (PDD) places the two phases on different GPU instances to eliminate cross-phase interference and specialize each instance for its phase [23, 38, 39, 58]. Prefill benefits from hierarchical dispatch and combine due to the large number of activations [15]. An activation crosses the scale-out network once to a remote node and is then
2.5
Prefill-Decode Colocation
Colocation avoids separate model replicas by sharing one instance across phases. To avoid interference, many techniques are proposed. Chunked prefill splits a long input into 3
MoE Server (rank=4)
Decode Server (rank=1)
MoE Server (rank=5)
Decode Server (rank=2)
MoE Server (rank=6)
Decode Server (rank=3)
Limitations. However, spatial partitioning cannot handle temporal resource demand variations along three dimensions. First, under EP, the number of activated experts and the number of tokens on a GPU change across layers, so the MoE load varies by rank and layer [20, 21, 47]. Second, attention and MoE have different resource demands even within the same layer [43, 60]. Third, each MoE module alternates among dispatch, expert computation, and combine, shifting its bottleneck between communication and computation [4]. A fixed partition cannot follow these changes. Changing an allocation requires CPU coordination, and kernel-completion waits make per-kernel changes impractical. Existing systems therefore repartition at the prefill layer boundary [10, 41]. As Figure 2 illustrates, a poor partition creates two failure modes. If prefill holds too many SMs when decode becomes ready, latency-sensitive decode waits behind a nonpreemptible prefill kernel. This head-of-line blocking can be severe because a prefill kernel may run for tens to hundreds of milliseconds while a decode kernel finishes in hundreds of microseconds, a difference of orders of magnitude. If resources remain reserved for decode when it has no ready work, prefill cannot use them, and the GPU develops a resource bubble. Rank-varying EP load and different module characteristics compound both problems. Existing solutions also leave communication unisolated. Prefill and decode dispatch and combine share the same NICs and links, where a large prefill transfer can throttle a small latency-sensitive decode transfer. Conventional twosided communication further requires matching two-sided progress. Independent scheduling can leave some MoE ranks serving prefill while others serve decode, so each phase waits for receiver work on ranks occupied by the other, and the system deadlocks. Finally, partitioning every GPU gives each phase fewer local resources. As Figure 2 shows, meeting the same latency target may then require a larger degree of parallelism, which increases communication and further amplifies network interference.
3
Prefill Server (rank=0)
Toke n1
Token 0
chunks that can be mixed with decode iterations to mitigate interference [2, 25]. However, it introduces memory access overhead from rereading prefix cache entries and model weights [58]. Recent systems colocate the phases through spatial GPU partitioning [10, 22, 41]. They commonly use Green Contexts [36] to reserve separate SMs for prefill and decode, avoiding duplicate model replicas while providing compute isolation.
Node 0
MoE Server (rank=7)
Prefill Server (rank=8)
Expert i Expert j
Decode Server (rank=9)
MoE Server (rank=13)
Decode Server (rank=10)
MoE Server (rank=14)
Decode Server (rank=11)
MoE Server (rank=15)
Node 1
Scale-up domain comm. Scale-out domain comm.
Figure 3. Architecture of ExpertPlex.
remaining GPUs to MoE servers. Prefill and decode servers in each node can be empty. Prefill and decode servers run their respective attention modules, while MoE servers host the MoE experts and execute expert computation for both phases. Because attention weights constitute only a small fraction of the model parameters while MoE weights account for over 95% of them, sharing MoE servers eliminates massive MoE weight duplication and multiplexes dynamically sparse expert computation. Disaggregating attention gives each phase full GPUs instead of GPU partitions, reducing its required degree of parallelism and communication while eliminating cross-phase attention network interference. This MoE-weight-independent boundary matches phase demands with fewer GPUs and enables finer-grained elastic scaling with a smaller failure blast radius. Each MoE server runs an Adaptive Persistent Kernel (APK) that keeps the GPU busy with adaptive scheduling (§4). APK preserves CUDA Graph compatibility and supports spatial and temporal multiplexing to maximize computation efficiency. Its tile-level scheduler follows real-time expert load, preempts prefill for decode within a sequence-lengthindependent bound, and reallocates idle SMs without CPU intervention while preserving prefill performance. ExpertPlex uses attention-initiated one-sided primitives for MoE communication (§5). Attention servers push dispatch data into and pull combine results from final MoE-side buffers. Direct final-buffer access removes MoE-side coordination and polling, avoids deadlock, and overlaps communication with computation across phases. To reduce network
ExpertPlex Overview
To address these problems, we propose ExpertPlex, a disaggregated serving system for MoE LLMs with adaptive persistent kernels. ExpertPlex shares MoE experts across phases while disaggregating attention by phase. As shown in Figure 3, each node assigns a configurable subset of GPUs to prefill servers, another subset to decode servers, and the 4
API InterCUDA ception [16, streams [35] 19, 49] CUDA Graph Compatibility Temporal Multiplexing Spatial Multiplexing Bounded Fast Preemption Bounded Fast Reallocation
✓
NVIDIA MPS [1]
Green Context [36]
NVIDIA MIG [37]
ExpertPlex
✓
✓
✓
✓
✓
Limited
✓ ✓ ✓ ✓ ✓
✓ ✓
Table 1. Comparison of ExpertPlex with existing GPU sharing solutions. interference, prefill scale-out traffic travels between attention servers whenever possible, while decode communicates directly between attention and MoE servers. Finally, a cross-stack optimizer models ExpertPlex from the tile level to the cluster level (§6). It jointly selects APK policies, placements, resource ratios, parallelism, and overlap strategies to maximize goodput under both phases’ SLOs.
4
Adaptive Persistent Kernel
4.1
GPU Sharing Design Space
streams with priorities [35] support graph replay and launchtime ordering, but provide no spatial isolation or bounded resource handoff once a kernel begins execution. Hardware partitioning provides spatial isolation but remains static. MPS [1] and Green Context [36] support graphcompatible spatial multiplexing, but fix the allocation during a kernel. They lack in-kernel time multiplexing and bounded preemption and reallocation. MPS also leaves cache and shared-memory interference on a shared SM unpredictable. MIG [37] provides hard but highly restricted spatial isolation. H100 exposes only 1g, 2g, 3g, 4g, and 7g profiles, so the only two-way split that lets two phases collectively use the GPU’s full compute capacity is 3g–4g. MIG lacks time multiplexing and preemption, and driver-level reconfiguration cannot bound reallocation latency. It therefore cannot follow the layer-level load changes described above for either phase.
Across MoE LLMs, prefill processes all input tokens, whereas each decode iteration processes only one new token per request. The duration gap between their operations therefore grows with input length. For example, with EP4 on MiniMaxM2.7, grouped GEMMs for decode with eight active experts take 17.7–34.7 𝜇s, while matched prefill GEMMs with 16K tokens take 1.8–2.9 ms, or 84–101× longer. The exact ratio is model dependent, but this trend is universal. Because every layer invokes an MoE operation, blocking can recur at every layer and accumulate into end-to-end decode latency. Static partitioning avoids this interference only by leaving resources idle whenever a phase is communicating or has no ready operation. A useful sharing mechanism must prevent repeated blocking without stranding this idle capacity. Efficient sharing therefore requires five complementary properties, summarized in Table 1. CUDA Graph compatibility avoids per-operation CPU intervention and its launch and synchronization overhead on short decode paths. Spatial multiplexing lets different SM subsets execute both phases concurrently. It lends SMs to the computing phase while the other communicates and changes the split across layers as routing shifts the expert load on each rank. Time multiplexing lets the same SM switch phases over time when its current phase has no ready operation. Bounded preemption prevents a long prefill operation from causing severe headof-line blocking, while bounded reallocation lets a newly ready phase reclaim SMs quickly enough to avoid bubbles. Launch-time mechanisms miss several properties. API interception [16, 19, 49] offers only launch-boundary time multiplexing. It lacks CUDA Graph compatibility, spatial multiplexing, and bounded preemption and reallocation. CUDA
4.2
Tile-Level Scheduling
These requirements favor a scheduling unit that is short and native to optimized GPU operations. APK chooses a tile, the smallest independently completable unit of these operations. Registers, shared memory, TMA buffers, and Tensor Core shapes bound its dimensions, so longer inputs add tiles without enlarging them. For example, our SM90 DeepGEMM configuration [56] uses CTA tiles of at most a 128×192 output block with a 128-wide reduction tile. Across MoE-side operations, tile boundaries occur every 2.2–25.3 𝜇s, independent of the total operation length (Figure 15). Input length stretches an operation by adding scheduling units rather than lengthening each scheduling interval. APK switches phases only after the current tile commits. No accumulator, TMA transaction, shared-memory buffer, or communication state remains live at this boundary, so switching needs no checkpoint, restore, or recomputation. Finer preemption would preserve pipeline state, while kernelor layer-level scheduling would retain long blocking intervals. APK keeps each operation’s native CTA or cluster shape, TMA multicast, and warp-specialized pipeline. The scheduler changes only the next tile, leaving each operation’s pipeline and synchronization intact throughout execution. Before APK starts, ExpertPlex builds an operation template including buffers and TMA descriptors for every MoE operation. At runtime, an attention server signals a ready 5
CTA Cluster w/ DSMEM CTA 0
TMA warp
Math warp
CTA i
…
TMA warp
SMEM Other Devices
P
Math warp
TMA warp
SMEM
pi
time plus one local cluster check epoch. This bound remains independent of the interrupted operation’s total length.
CTA i+1
Math warp
5 Attention-initiated MoE Communication
SMEM
5.1
Global Memory System scope
Device scope
Disaggregating prefill and decode attention onto different servers eliminates cross-phase interference for attention communication. The remaining contention point is therefore MoE communication, because both phases dispatch routed activations to, and combine expert outputs from, the same MoE server pool. Existing MoE communication usually uses a two-sided protocol for two reasons. Token routing reveals destinations and transfer sizes only at runtime, and the communication operator is decoupled from the compute operator that consumes its output. The sender therefore streams data into a bounded receiver-side ring buffer, while the receiver polls and drains that buffer, scatters arrivals into the final tensor, and returns credits that prevent the sender from overwriting unread slots. The ring buffer bounds memory use, but progress now requires both sides. With concurrent prefill and decode, independent GPU scheduling can leave some MoE ranks running prefill kernels and others running decode kernels. Each phase then waits for receiver kernels on ranks occupied by the other phase, so neither ring buffer is drained nor returns the credits needed for its senders to finish. The blocked sender cannot retire and release its GPU resources for the other phase’s receiver kernel. This cross-phase cycle can deadlock an otherwise correct protocol. Reserving receiver-side polling SMs avoids this liveness failure only by wasting compute capacity. MoE traffic is bursty, so these SMs spin between dispatches and cannot execute ready expert tiles. Letting compute borrow them is unsafe because receiver progress would again depend on whether enough SMs become available at every rank. Network contention remains as well. Optimized MoE transfers can saturate the fabric [15], while scale-out bandwidth is scarce. For example, DeepSeek-V3 reports 160 GB/s intranode NVLink bandwidth but only 50 GB/s cross-node InfiniBand bandwidth on H800, a 3.2× bandwidth gap [13]. A prefill burst can delay latency-sensitive decode even when the phases use different GPUs.
Figure 4. Tile-level preemption mechanism. operation. CTA 0 instantiates its template into the decode or prefill queue and updates runtime fields such as the layer, token count, and SM budget. APK schedules at CTA-cluster granularity because CTAs connected by TMA multicast must execute the same tile. The cluster leader selects a tile and broadcasts its metadata to peer CTAs. Different clusters execute different phases for spatial multiplexing, while each cluster switches phases at tile boundaries for time multiplexing and fast reallocation. Because APK is long-lived and all scheduling occurs on the GPU, this process requires neither CPU intervention nor kernel relaunch and remains compatible with CUDA Graph replay throughout execution. 4.3
Concurrent Communication over a Shared Network
Bounded Tile-level Preemption
Independent tile-boundary checks are unsafe because highperformance operations pipeline both warps and CTAs across tiles. Within a CTA, a TMA warp may load tile 𝑘 + 1 while math warps consume tile 𝑘. If one warp switches phases, another can wait forever for data or an mbarrier event from the old phase. Across a CTA cluster, an independently switching CTA can similarly strand peers at a TMA multicast or cluster barrier. Stalling the entire cluster after every tile prevents both deadlocks but serializes a microsecond-scale pipeline and defeats fine-grained scheduling. APK instead propagates one cooperative decision through the memory hierarchy shown in Figure 4. An attention server signals an urgent decode operation through the system-scope word 𝑃. CTA 0 checks 𝑃 at a tile boundary of its current operation and writes one device-scope word 𝑝𝑖 for each cluster 𝑖. Each cluster reads only its 𝑝𝑖 , avoiding repeated systemscope accesses and a grid-wide barrier. At its next tile boundary, the leader CTA broadcasts the decision through DSMEM so every CTA participating in TMA multicast stays on the same operation. Within each CTA, the first pipeline warp reads the decision, broadcasts it within the warp, and uses an mbarrier handoff to notify later warps before they claim another tile. The decision remains fixed within a check epoch, so no CTA can claim another old-phase tile after observing it. The slowest CTA therefore finishes at most its current tile before the cluster converges on the new phase. The propagation follows the GPU memory hierarchy through system memory, device memory, DSMEM, shared memory, and warp broadcast. It prevents mixed-phase dependencies both within a CTA and across a CTA cluster without draining the pipeline after every tile. Preemption is bounded by one tile execution
5.2
One-Sided Dispatch and Combine
Our key insight is to remove MoE-side progress with onesided transfers initiated by attention servers. ExpertPlex uses both push and pull primitives of the scale-up links, i.e., NVLink, and scale-out links, i.e., RDMA. APK already preallocates each operation’s buffers for its maximum routed-token volume. This requirement follows from APK’s persistent execution because its operation templates, tensor addresses, and communication descriptors must be fixed before runtime scheduling. ExpertPlex exposes these final dispatch and 6
Decode Server’s Comm. Stream
combine buffers to attention servers and removes the intermediate ring buffer. Transfers consequently need neither receiver-side draining and scattering nor credit feedback to throttle a sender. MoE servers expose only buffers and readiness words, with no matching communication kernel or reserved polling SMs. Dispatch uses an attention-side push. After routing determines each expert’s destination and final offset, the attention server writes activations into that location through NVLink peer stores or one-sided RDMA writes. It publishes a ready signal only after the payload is visible. APK does not coordinate this transfer and schedules the MoE task after observing that signal at a tile boundary. Combine reverses the data flow but remains attentioninitiated. After expert computation writes the final combine buffer and publishes a done signal, a small WaitDone kernel on the attention GPU observes completion, and the attention server pulls the result through NVLink loads or one-sided RDMA reads. WaitDone is a single-thread kernel, so it can coexist with TBO or SBO and preserve their attention-side overlap. It runs on the communication stream and blocks only the dependent pull, while the compute stream continues the work exposed by the selected overlap strategy. Meanwhile, removing MoE-side communication kernels lets APK overlap one phase’s MoE computation with the other phase’s dispatch or combine, as Figure 5 shows. For example, APK can execute prefill tiles while decode waits to combine, or execute urgent decode tiles while a prefill transfer is in flight. This mechanism is compatible with each phase’s original overlap strategy but adds cross-phase opportunities. Prefill and decode follow independent dependency chains, so this overlap is not limited by the intra-phase dependencies that constrain TBO and SBO. It can keep the SMs busy as long as there is work available. 5.3
Combine
Dispatch
Combine
WaitDone
WaitDone
Read done signal
Write start signal Read done signal
Data Transfer
MoE Server’s GPU
Data Transfer
Preempt!
…
CTA Clusters
Data Transfer
Read done signal
Write start signal
Dispatch Prefill Server’s Comm. Stream
WaitDone
Decode’s MoE computation
Data Transfer Combine
Prefill’s MoE computation
Figure 5. Cross-phase overlap enabled by attention-initiated one-sided transfers. decode. Topology thus separates the phases whenever possible, while virtual-lane priority protects latency-sensitive decode when they must share a scale-out link.
6
Cross-stack Placement Optimizer
6.1
Cross-stack Scheduling Problem
Although the mechanisms above expose enough control to adaptively schedule prefill and decode with stricter isolation at finer granularity, they also couple decisions conventionally made by different layers of the serving stack. For example, changing the number of prefill servers simultaneously changes the degree of parallelism, the KV-cache capacity, and the communication path. Optimizing placement, overlap, or GPU sharing in isolation can therefore select a locally efficient configuration that violates an SLO once these effects interact. We instead jointly search the attention and MoE parallelism, server placement, overlap strategy, and APK sharing policy. We formalize a configuration as a layout ℓ and a decode SM budget 𝑞 under contention. For (ℓ, 𝑞), let 𝐵𝑝 and 𝐵𝑑 denote the largest prefill and decode batches whose modeled iteration latencies 𝑇𝑝 and 𝑇𝑑 meet their respective SLOs, and let 𝑂¯ be the average output length. The request-level goodput is defined as 𝐵𝑝 𝐵𝑑 𝐺 (ℓ, 𝑞) = min , . (1) 𝑇𝑝 𝑇𝑑 𝑂¯ The minimum captures pipeline balance because each request consumes one prefill iteration but 𝑂¯ decode iterations. We admit only candidates satisfying both phase SLOs and maximize 𝐺 over the feasible candidates, rather than overprovisioning either phase in isolation.
Hierarchical Prefill Path and Traffic Isolation
One-sided transfers remove MoE-side coordination, but prefill and decode can still contend for scarce RDMA bandwidth. Our key insight is to route most prefill scale-out traffic between prefill attention servers, reserving direct attention-toMoE paths primarily for decode. Latency-sensitive decode accesses each MoE server directly through local NVLink or remote RDMA. For a remote node with a prefill server, the source sends each activation once over RDMA to the prefill server with the same local rank on that node. The receiver multicasts it out over NVLink to all activated experts chunk by chunk. For combine, the same server gathers outputs over NVLink and returns the combined result through the prefill-only RDMA path. This path preserves hierarchical deduplication because activations for multiple experts on one remote node cross the scarce scale-out link only once. The placement optimizer spreads prefill servers across nodes to make this path common (§6). When a destination node lacks a prefill server, ExpertPlex uses direct one-sided RDMA but assigns prefill a lower-priority InfiniBand virtual lane than
6.2
Tile-aware Latency Model
Evaluating every candidate on a full cluster would make the search prohibitively expensive. ExpertPlex instead profiles attention and MoE kernels, including APK overhead, on only a few representative GPUs and interpolates their per-rank latency. Network time requires no profiling GPU. The model computes it directly from routed bytes, measured RDMA or 7
Execution time (us)
GateUp #expert=1 Down #expert=1 GateUp #expert=2 Down #expert=2
GateUp #expert=4 Down #expert=4 GateUp #expert=8 Down #expert=8
GateUp #expert=16 Down #expert=16 GateUp #expert=32
Down #expert=32 GateUp #expert=64 Down #expert=64
Algorithm 1 Cross-stack placement search 1: 𝑏𝑒𝑠𝑡 ← ⊥ 2: for all ℓ ∈ L do
200 150 100
3: 4:
50
5:
0
1
2
4
8
16
Per expert batch size
32
64
6: 7: 8: 9: 10: 11:
128
Figure 6. Latency of MoE Grouped GEMMs under different numbers of active experts and tokens. NVLink bandwidth, and the hierarchical path in §5. For each layout, it derives the work on every rank from the attention and MoE parallelism, placement, routing, and sharing policy. For a profiled component 𝑐, we fit a small set of input samples with 𝑡ˆ𝑐 (𝑥, 𝑠) = 𝛼𝑐 + 𝛽𝑐 𝑥 + 𝛾𝑐 𝑥𝑠 + 𝛿𝑐 𝑥𝑠 2 .
12: return 𝑏𝑒𝑠𝑡
under 𝑞 SMs, and the communication latency. The overall iteration latency is then estimated as the sum of latencies on the critical path. Under TBO, ExpertPlex overlaps one microbatch’s attention with another’s communication and MoE computation, so the iteration latency is the maximum of these two paths. Under SBO, ExpertPlex overlaps sharedexpert computation with routed-expert computation and communication within one microbatch, so the iteration latency is their maximum plus the attention latency. Lines 8–11 retain only configurations feasible for both phases and select the one with maximum 𝐺. The selected configuration jointly fixes the parallelism, placement, overlap strategy, and the expected decode SM budget.
(2)
For attention, 𝑥 is the local batch size and 𝑠 is the sequence length. Our key observation is that raw routed-token count is insufficient for MoE computation because its latency follows the number of executed tiles. If expert 𝑒 receives 𝑚𝑒 token rows and the kernel tile height is 𝑀𝑡 , the MoE footprint is ∑︁ 𝑚𝑒 𝑥 moe = (3) 𝑀𝑡 𝑒 |𝑚𝑒 >0
and we evaluate the fitted form with 𝑥 = 𝑥 moe and 𝑠 = 1. Figure 6 illustrates this distinction. Batches with the same number of tokens can have different latency when they activate different numbers of experts. This behavior is fundamental to tiled MoE kernels rather than specific to one model or GPU. Tensor Core and TMA pipelines operate on aligned tiles, and each active expert triggers at least one tile’s metadata, weight transfer, and computation even if it receives only a few tokens. Thus sparse routing changes latency through both token volume and its per-expert packing. By predicting the per-expert token vector and interpolating on 𝑥 moe , the model captures this general effect with a small GPU profile set instead of profiling every full-cluster layout. 6.3
if ¬FitsMemory(ℓ) then continue for all 𝑞 ∈ Qℓ do (𝐵𝑝 , ok𝑝 ) ← MaxFeasibleBatch(ℓ, 𝑞, prefill) (𝐵𝑑 , ok𝑑 ) ← MaxFeasibleBatch(ℓ, 𝑞, decode) if ok𝑝 ∧ ok𝑑 then 𝑔 ← 𝐺 (ℓ, 𝑞) if 𝑏𝑒𝑠𝑡 = ⊥ ∨ 𝑔 > 𝑏𝑒𝑠𝑡 .𝑔 then 𝑏𝑒𝑠𝑡 ← (ℓ, 𝑞, 𝐵𝑝 , 𝐵𝑑 , 𝑔)
6.4
Online SM Reallocation
The offline solution targets an average workload, whereas routing changes 𝑥 moe across layers and servers, and either phase may temporarily have no ready MoE computation. APK therefore interprets 𝑞 as a contention policy rather than a static partition. When only one phase is ready, it uses all CTA clusters. Under contention, it scales 𝑞 by the ratio of the current decode footprint 𝑥 moe to the offline expectation ★ as follows. 𝑥 moe 𝑞𝑥 moe 𝑞 ′ = min 𝑄 max, ★ , (4) 𝑥 moe 𝑐
Offline Search Algorithm
Here, ⌈·⌉𝑐 rounds up to a CTA-cluster multiple, 𝑄 max preserves prefill progress, and prefill receives the remaining clusters. The reason for protecting decode share first is that decode is more latency-sensitive, and that outputs for alreadystarted requests should be served before inputs for new requests. For the prefill phase, APK uses all the remaining SMs to preserve throughput as high as possible.
Algorithm 1 jointly searches these choices for the workload’s expected input and output lengths. The generator enumerates layouts L, where each ℓ specifies the number of MoE servers, the number of attention servers, the prefill-to-decode server ratio, and the layout of all servers. FitsMemory first rejects layouts that cannot hold attention and expert weights or sufficient KV caches. For every remaining layout, Qℓ enumerates decode SM budgets. For each (ℓ, 𝑞), MaxFeasibleBatch binary-searches up to the configured and KV-cache bounds for each phase (lines 6–7). At each probe, the model from §6.2 estimates the attention latency, the tile-aware MoE computation latency
7
Evaluation
7.1
Experiment Setup
Implementation. We implement ExpertPlex by modifying DeepGEMM [56], DeepEP v1 [15], and SGLang [57]. The 8
90% SLO
80% 60% 40% 20% 0%
0
5
10
15
20
25
30
ExpertPlex (1P3D4E)
100%
SGLang-PDD (1P1D x DP4EP4) 90% SLO
80% 60% 40% 20% 0%
0
5
10
Request Rate (req/s/node)
15
20
25
SGLang-PDMux (TP8EP8)
100%
SLO Attainment (%)
SGLang-Colocated (DP8EP8)
TPOT SLO Attainment (%)
TTFT SLO Attainment (%)
SGLang-ChunkedPrefill (Chunk=2048, DP8EP8)
100%
60% 40% 20% 0%
30
90% SLO
80%
0
5
10
Request Rate (req/s/node)
15
20
25
30
Request Rate (req/s/node)
Figure 7. SLO attainment of ExpertPlex and baselines on the MiniMax-M2.7 and ShareGPT datasets. SGLang-Colocated (8DP8EP)
90% SLO
80% 60% 40% 20% 0% 0.0
0.5
1.0
1.5
2.0
2.5
ExpertPlex (3P1D4E)
100%
SGLang-PDD (1P1D x DP4EP4) 90% SLO
80% 60% 40% 20% 0% 0.0
0.5
Request Rate (req/s/node)
1.0
1.5
2.0
SGLang-PDMux (TP8EP8)
100%
SLO Attainment (%)
100%
TPOT SLO Attainment (%)
TTFT SLO Attainment (%)
SGLang-ChunkedPrefill (Chunk=2048, DP8EP8)
90% SLO
80% 60% 40% 20% 0% 0.0
2.5
0.5
Request Rate (req/s/node)
1.0
1.5
2.0
2.5
Request Rate (req/s/node)
Figure 8. SLO attainment of ExpertPlex and baselines on the MiniMax-M2.7 and LooGLE datasets.
80% 60% 40% 20% 0%
0
1
2
3
4
5
SGLang-Colocated (DP16EP16)
ExpertPlex (3P6D15E)
100%
90% SLO
80% 60% 40% 20% 0%
0
Request Rate (req/s/node)
1
2
3
4
SGLang-PDMux (TP16EP16)
100%
SLO Attainment (%)
90% SLO
TPOT SLO Attainment (%)
TTFT SLO Attainment (%)
SGLang-ChunkedPrefill (Chunk=4096, DP16EP16)
100%
60% 40% 20% 0%
5
90% SLO
80%
0
Request Rate (req/s/node)
1
2
3
4
5
Request Rate (req/s/node)
Figure 9. SLO attainment of ExpertPlex and baselines on the GLM-5.1-FP8 and ShareGPT datasets.
60% 40% 20% 0% 0.00
0.05
0.10
0.15
0.20
Request Rate (req/s/node)
0.25
SGLang-Colocated (16DP16EP)
ExpertPlex (6P3D15EP)
100%
90% SLO
80% 60% 40% 20% 0% 0.00
0.05
0.10
0.15
0.20
Request Rate (req/s/node)
0.25
SGLang-PDMux (TP16EP16)
100%
SLO Attainment (%)
90% SLO
80%
TPOT SLO Attainment (%)
TTFT SLO Attainment (%)
SGLang-ChunkedPrefill (Chunk=4096, 16DP16EP)
100%
90% SLO
80% 60% 40% 20% 0% 0.00
0.05
0.10
0.15
0.20
0.25
Request Rate (req/s/node)
Figure 10. SLO attainment of ExpertPlex and baselines on the GLM-5.1-FP8 and LooGLE datasets. adaptive persistent kernel (APK) is implemented based on ibgda_get_nbi_warp to support fine-grained pulls from reDeepGEMM and DeepEP v1. We fuse two grouped GEMMs, mote combine buffers over a specific RDMA connection. In activation functions, MoE data preprocessing, MoE data postSGLang, we add support for the disaggregated attention and processing and the CTA-cluster scheduler into one long-lived shared-MoE layout, the MoE token dispatcher for this layout, persistent kernel on each MoE GPU. For each operation, the and a OperationsStrategy for communication overlap. host determines tile sizes before launch and JIT-compiles them as constants at startup to reduce register pressure. Models. We use MiniMax-M2.7 for single-node experiments We also replace DeepEP’s receiver-driven all-to-all path and GLM-5.1-FP8 for multi-node experiments. Both models with the attention-initiated M-to-N path described in Figure 3 use 256 routed experts per MoE layer and activate eight and Figure 5. Dispatch writes routed activations and metarouted experts per token. MiniMax-M2.7 has a 230 GB FP8 data directly into the target MoE buffers without receivermodel footprint in our deployment, and each token activates side coordination. Combine is initiated from the attention about 7.0B routed expert parameters across layers. GLMside through pull semantics, so the MoE server only publishes 5.1-FP8 has a 756 GB FP8 model footprint and 724.8B routed finished expert outputs and ready signals. For the InfiniBand expert parameters, of which each token activates about 22.6B. GPUDirect Async (IBGDA) path, we customize nvshmemi_ The former uses full attention; the latter uses DSA. 9
8000 6000
Slowdown vs exclusive (x)
Latency of prefill's grouped GEMM (us)
Exclusive ExpertPlex
Exclusive Green Context ExpertPlex CUDA Stream MPS
10000
4000 2000 00
200
400
600
800
1000
1200
Latency of decode's grouped GEMM (us) Figure 11. Pareto frontier of GPU sharing mechanisms.
Testbed. Single-node experiments run on one NVIDIA H800 node with eight GPUs connected by NVLink. Multi-node experiments use up to three machines. Each machine has eight NVIDIA H800 GPUs connected by NVLink. Machines are connected by eight 200 Gbps InfiniBand NICs per node. All throughput numbers are reported as requests per second per node when systems use different numbers of nodes.
15.0 12.5 10.0 7.5 5.0 2.5 0.0
MPS CUDA Stream
Green Context
13.79x
4.07x
3.33x 1.00x 1.08x 1.24x
1.34x
1.00x 1.12x
1.04x
Prefill's grouped GEMM Decode's grouped GEMM Figure 12. Case study of preemption slowdown.
TTFT SLO is 1 s and the TPOT SLO is 50 ms. For MiniMaxM2.7 on LooGLE, the TTFT SLO is 10 s and the TPOT SLO is 100 ms. For GLM-5.1-FP8 on ShareGPT, the TTFT SLO is 2 s and the TPOT SLO is 100 ms. For GLM-5.1-FP8 on LooGLE, the TTFT SLO is 20 s and the TPOT SLO is 100 ms. These SLOs follow the range used by prior disaggregated and high-goodput serving work [10, 43, 55, 60].
Baselines. Because ExpertPlex is implemented in SGLang, all baselines use SGLang-based implementations to isolate system-level differences. SGLang-Colocated is the original SGLang serving mode. It enables DP for attention and EP for MoE. When memory allows, it enables two-batch overlap to improve throughput. SGLang-ChunkedPrefill additionally enables chunked prefill with the best chunk size found in prior tuning for the corresponding workload and model. Other settings match SGLang-Colocated. SGLang-PDD implements widely used prefill-decode disaggregation. Because MoE weights consume most GPU memory, MiniMax-M2.7 uses a 1P1D deployment. GLM-5.1-FP8 runs out of memory under this PDD layout on 24 GPUs, so we do not report GLM-5.1-FP8 PDD numbers. SGLang-PDMux is based on the open-source MuxWise implementation. MuxWise targets dense models and uses Green Context to partition GPU resources, so we modify it to support MoE models. Because its implementation is compatible only with tensor-parallel attention, it uses TP for attention and EP for MoE. Several baselines cannot express ExpertPlex’s 24-GPU fine-grained layout under their parallelism constraints. For GLM-5.1-FP8, those baselines run on the largest compatible 16-GPU layout, and we compare request rates per node for a fair comparison.
7.2
End-to-End Performance
Figures 7, 8, 9, and 10 show the goodput of all systems on different models and workloads. Across the four settings, ExpertPlex improves goodput by matching phase resources at the MoE-attention boundary instead of duplicating the whole model or partitioning every GPU by phase. For MiniMax-M2.7 on ShareGPT, ExpertPlex reaches 11.3 requests per second per node under the joint SLO. This is 5.65× higher than SGLang-ChunkedPrefill, 2.72× higher than SGLang-Colocated, 2.01× higher than SGLang-PDD, and 1.41× higher than SGLang-PDMux. Chunked prefill performs worst. Each chunk rereads model weights and the KV cache and repeats MoE communication. Large chunks still interfere with decode, whereas small chunks reduce prefill efficiency. Colocated serving runs prefill aggressively to build larger batches and improve throughput, but lets long prefill delay decode iterations. SGLang-PDD removes prefill-decode interference, but it duplicates massive MoE weights and leaves less room for KV cache. Instancelevel disaggregation also leads to resource mismatch, so the prefill phase uses only part of the available resources even when decode is not the bottleneck. SGLang-PDMux improves over colocated serving by using Green Context, yet its GPU partitions cannot follow the actual number of active experts on each MoE rank. It also introduces cross-phase communication interference and requires a larger degree of parallelism. ExpertPlex avoids these failure modes by sharing the MoE pool with APKs while disaggregating attention modules. On MiniMax-M2.7 with LooGLE, longer requests magnify the difference. SGLang-ChunkedPrefill cannot sustain the SLO over the evaluated load range because longer inputs amplify its memory-access overhead. They also amplify interference in SGLang-Colocated, over which ExpertPlex
Workloads. Following prior works, we sample input and output lengths from ShareGPT and LooGLE to represent short and long requests. Request arrivals follow a Poisson process. Because SGLang-PDD leaves less GPU memory for the KV cache after duplicating MoE weights across prefill and decode instances, we cap sampled sequence lengths at the PDD KV-cache capacity. Metrics. The main metric is P90 goodput, defined as the highest arrival rate at which at least 90% of requests meet both their time to first token (TTFT) and time per output token (TPOT) SLOs. For MiniMax-M2.7 on ShareGPT, the 10
ExpertPlex-GateUp ExpertPlex-DeepGEMM
Execution time (us)
25000 20000 15000 10000 5000 00
10000
20000
30000
Per expert batch size
(a) Contiguous layout.
175 150 125 100 75 50 25 00
Dispatch
GateUp Overhead Down Overhead
10
15
Number of experts
DeepEPv1 800
ExpertPlex
600
6000
400
4000 200
2000 8192
16384 32768
Number of tokens
(a) Normal mode.
0
16
64
256
Number of tokens
1024
(b) Low-latency mode.
Figure 14. Overhead of attention-initiated communication.
(b) Masked layout.
Figure 13. Overhead of adaptive persistent kernels.
keep latency of the decode phase close to exclusive execution, but their fixed allocations slow prefill by 3.33× and 4.07×. In contrast, ExpertPlex adds only 8% overhead to the decode phase and slows prefill by only 1.12×. The reason is that decode GEMMs are short and intermittent. APK gives all idle CTA clusters to prefill, only preempts prefill when decode arrives, and reallocates idle SMs back to prefill when decode completes the execution.
improves goodput by 4.12×; the improvement over SGLangPDMux is 1.28× on this workload. GLM-5.1-FP8 shows similar improvements in multi-node deployments. Compared with SGLang-ChunkedPrefill, ExpertPlex improves goodput by 3.3× on ShareGPT and 5.0× on LooGLE. Compared with SGLang-Colocated, ExpertPlex improves goodput by 1.5× on ShareGPT and 2.5× on LooGLE. On ShareGPT, ExpertPlex and SGLang-PDMux achieve similar joint goodput of about 1.5 requests per second per node. PDMux’s tensor-parallel attention gives each short request more parallelism and thus a TTFT advantage, but adds communication and allocates resources without regard to MoE sparsity. On LooGLE, this advantage fades and ExpertPlex improves goodput over SGLang-PDMux by 1.66×. Long requests expose network interference and Green Context’s inability to maximize GPU utilization while preserving TPOT. 7.3
8000
0 4096
5
Combine
10000
Execution time (us)
DeepGEMM-GateUp DeepGEMM-Down
7.4
Overhead of Tile-Level Scheduling
Tile-level scheduling adds queue checks, task selection, and CTA-budget enforcement to grouped GEMM. We measure its overhead against DeepGEMM using the configuration from §7.3. Figure 13(a) evaluates the contiguous layout used by prefill. The x-axis varies per-expert batch size, the dominant factor in prefill GEMM time. Scheduling adds less than 12% overhead because it runs once per tile group without changing the TMA or Tensor Core pipelines. Figure 13(b) evaluates the masked layout used by decode. For decode, execution time depends more on the number of active experts than on the number of tokens. Scheduling adds less than 20 𝜇s across all measured active-expert counts. This small cost replaces a much longer wait behind prefill kernels. Moreover, because each token can activate eight experts, its relative overhead falls below 10% for grouped GEMMs with many active experts.
Effectiveness of APKs
We evaluate APK’s GPU sharing mechanism with concurrent prefill- and decode-shaped grouped GEMMs from GLM-5.1FP8 on one GPU. The decode grouped GEMM is launched 10 𝜇s after the prefill grouped GEMM. The number of tokens is 128 in the decode phase and 8192 in the prefill phase, and all activate eight experts. We compare APK with exclusive execution, priority CUDA streams, MPS, and Green Context. Figure 11 plots the Pareto frontier. CUDA stream priorities order kernel launches but provide no isolation after launch, so decode waits behind a running prefill GEMM, leading to significant head-of-line blocking. Green Context spatially isolates the phases, but its fixed partition prevents either phase from borrowing the other’s idle SMs, leading to resource bubbles. MPS multiplexes space more flexibly, but still cannot time-multiplex an SM at tile granularity within a running kernel. APK is the only mechanism that reaches the low-latency region for the decode phase while keeping the prefill phase high-performing. Figure 12 shows the result under the decode-latency protection target. CUDA streams increase decode latency by 13.79× over exclusive execution. MPS and Green Context
7.5
Overhead of Attention-Initiated MoE Communication
Figure 14 measures the attention-initiated MoE communication described in §5. On 16 GPUs, we compare DeepEP v1 with ExpertPlex’s attention-initiated one-sided dispatch and combine under both normal and low-latency execution. In normal mode, dispatch and combine track DeepEP v1 within about 5%, preserving hierarchical communication’s bandwidth benefit. Low-latency dispatch and combine expose measurement noise, yet the difference remains within about 45 𝜇s. Thus, removing MoE-side coordination and using pull-based combine integrates communication with APK without sacrificing communication efficiency, while leaving MoE-server SMs available to the other phase. 11
Preemption Interval (us)
Bless: 700 us
750 500 250 100
these designs are built on instance-level PDD, so they inherit the same limitations and further introduce pipeline bubbles [4]. In contrast, ExpertPlex shares the MoE servers between phases, which maximizes GPU occupancy and reduces bubbles. Its attention-initiated communication further introduces communication-computation overlap across phases. Attention-side techniques such as sequence parallelism [48, 50] and attention offloading [29, 30, 46] are orthogonal to ExpertPlex. MoE load-balancing [20, 21, 28, 47, 52] solutions are also orthogonal, while ExpertPlex can absorb phase-level load shifts through multiplexing.
LithOS: 250 us GPREEMPT: 100 us
80 60 40
EP Scatter Stage 1 EP Scatter Stage 2 UpGate GEMM
SiLU + Mul + Quant Down GEMM EP Gather
REEF: 35 us
20 0 128
256
512
Per expert batch size
1024
2048
Figure 15. Preemption interval of different tasks. 7.6
Kernel scheduling. Many systems improve GPU efficiency by scheduling work below the operator boundary. Megakernel systems break kernel boundaries, fuse small or memorybound operations, and exploit inter-operator parallelism [11, 18, 24, 32, 42, 45, 59]. These techniques target small batch sizes or intra-phase efficiency, but ExpertPlex targets highgoodput MoE serving with concurrent prefill and decode phases, where the scheduler must protect decode SLOs while keeping prefill throughput high. Other systems overlap computation and communication through fine-grained pipelining or kernel fusion [4, 9, 54]. They overlap operations within one phase or one layer pipeline. ExpertPlex keeps these optimizations compatible, but it can also overlap communication from one phase with computation from the other phase. GPU sharing and preemption systems provide another line of fine-grained scheduling [1, 7, 12, 16, 19, 37, 49]. These systems target multitasking across models or applications rather than multiple phases of one MoE model. They also do not support some modern GPU kernel features such as TMA multicast, CTA clusters, warp specialization, and CUDA Graph. More importantly, their preemption mechanisms also introduce checkpoint, restore, or recomputation overhead. ExpertPlex uses tile-level scheduling to avoid these issues to guarantee bounded low preemption and reallocation overhead with temporal and spatial multiplexing support.
Analysis of Preemption Intervals
ExpertPlex tunes GEMM and data-processing tiles for operation performance rather than shrinking them to improve preemption. Under these tile sizes, Figure 15 shows that all MiniMax-M2.7 MoE operations have intervals below 25.3 𝜇s, and particularly GEMM intervals stay below 10.7 𝜇s. The interval depends on tile execution time, not token count, allowing a short decode operation to preempt work from a prefill request with thousands of input tokens. We compare against the best reported delays of LithOS, Bless, GPreempt, PipeSwitch, and REEF [7, 12, 16, 19, 53] only as reference points, since they do not target MoE LLM workloads and lack the support of modern MoE kernel features such as TMA multicast, CTA clusters, warp specialization, and CUDA Graph. APK instead schedules at natural tile boundaries, performs useful work throughout each interval, and requires no checkpoint, restore, or recomputation. Even REEF’s best reported delay is 35 𝜇s and requires recomputing the preempted kernel, which is costly when millisecondscale prefill GEMMs are frequently interrupted by decode. APK therefore achieves finer-grained preemption with lower pause and resume costs for MoE serving workloads.
8
Related Work
LLM serving systems. Many techniques are proposed to adapt to accelerate LLM serving. Prefill-decode disaggregation (PDD) provisions independent prefill and decode instances to avoid interference between two phases [23, 38, 39, 58]. As discussed in §2, instance-level PDD is inefficient to deploy on small clusters, enlarges the rank-failure blast radius and only supports coarse elastic scaling. Chunked prefill instead splits prefill into smaller chunks to reduce interference [2, 25], while repeated chunk execution introduces extra memory access. Recent GPU partitioning systems let prefill and decode share a GPU by assigning resources through mechanisms such as Green Context [10, 22, 31, 41]. However, as described in §2, these systems can introduce bubbles, head-of-line blocking, network interference, and a larger degree of parallelism. More recent attention-expert disaggregation systems disaggregate attention and expert computation to fit their different characteristics [43, 55, 60]. However,
9
Conclusion
We presented ExpertPlex, a high-goodput MoE LLM serving system that disaggregates attention by phase while sharing massive MoE experts. ExpertPlex shares MoE weights for memory efficiency, disaggregates attention to reduce attention communication of wide parallelism, uses adaptive persistent kernels to multiplex dynamic sparse MoE computation with performance isolation, and initiates one-sided MoE communication from attention servers to avoid crossphase interference and overlap communication with computation. Evaluations on MiniMax-M2.7 and GLM-5.1-FP8 show that ExpertPlex improves goodput by up to 5.65× over chunked prefill, 2.01× over instance-level prefill-decode disaggregation, 1.66× over Green Context-based prefill-decode colocation, and 4.12× over naive colocated serving systems. 12
References
[20] Jiaao He, Jiezhong Qiu, Aohan Zeng, Zhilin Yang, Jidong Zhai, and Jie Tang. 2021. Fastmoe: A fast mixture-of-expert training system. arXiv (2021). [21] Jiaao He, Jidong Zhai, Tiago Antunes, Haojie Wang, Fuwen Luo, Shangfeng Shi, and Qin Li. 2022. Fastermoe: modeling and optimizing training of large-scale dynamic pre-trained models. In ACM PPoPP. [22] Ke Hong, Lufang Chen, Zhong Wang, Xiuhong Li, Qiuli Mao, Jianping Ma, Chao Xiong, Guanyu Wu, Buhe Han, Guohao Dai, Yun Liang, and Yu Wang. 2025. semi-PD: Towards Efficient LLM Serving via Phase-Wise Disaggregated Computation and Unified Storage. arXiv (2025). [23] Cunchen Hu, Heyang Huang, Liangliang Xu, Xusheng Chen, Jiang Xu, Shuang Chen, Hao Feng, Chenxi Wang, Sa Wang, Yungang Bao, Ninghui Sun, and Yizhou Shan. 2024. Inference without Interference: Disaggregate LLM Inference for Mixed Downstream Workloads. In arXiv. [24] Hongyi Jin, Bohan Hou, Guanjie Wang, Ruihang Lai, Jinqi Chen, Zihao Ye, Yaxing Cai, Yixin Dong, Xinhao Cheng, Zhihao Zhang, Yilong Zhao, Yingyi Huang, Lijie Yang, Jinchen Jiang, Gabriele Oliaro, Jianan Ji, Xupeng Miao, Vinod Grover, Todd C. Mowry, Zhihao Jia, and Tianqi Chen. 2026. Event Tensor: A Unified Abstraction for Compiling Dynamic Megakernel. In Conference on Machine Learning and Systems. [25] Aditya K. Kamath, Ramya Prabhu, Jayashree Mohan, Simon Peter, Ramachandran Ramjee, and Ashish Panwar. 2024. POD-Attention: Unlocking Full Prefill-Decode Overlap for Faster LLM Inference. In arXiv. [26] Jared Kaplan, Sam McCandlish, Tom Henighan, Tom B. Brown, Benjamin Chess, Rewon Child, Scott Gray, Alec Radford, Jeffrey Wu, and Dario Amodei. 2020. Scaling Laws for Neural Language Models. (2020). [27] Kimi Team. 2025. Kimi K2: Open Agentic Intelligence. In arXiv. [28] Jiamin Li, Yimin Jiang, Yibo Zhu, Cong Wang, and Hong Xu. 2023. Accelerating Distributed MoE Training and Inference with Lina. In USENIX ATC. [29] Yunkai Liang, Zhangyu Chen, Pengfei Zuo, Zhi Zhou, Xu Chen, and Zhou Yu. 2025. Injecting Adrenaline into LLM Serving: Boosting Resource Utilization and Throughput via Attention Disaggregation. arXiv (2025). [30] Bin Lin, Tao Peng, Chen Zhang, Minmin Sun, Lanbo Li, Hanyu Zhao, Wencong Xiao, Qi Xu, Xiafei Qiu, Shen Li, Zhigang Ji, Yong Li, and Wei Lin. 2024. Infinite-LLM: Efficient LLM Service for Long Context with DistAttention and Distributed KVCache. In arXiv. [31] Zejia Lin, Hongxin Xu, Guanyi Chen, Xianwei Zhang, and Yutong Lu. 2025. Bullet: Boosting GPU Utilization for LLM Serving via Dynamic Spatial-Temporal Orchestration. arXiv (2025). [32] Lingxiao Ma, Zhiqiang Xie, Zhi Yang, Jilong Xue, Youshan Miao, Wei Cui, Wenxiang Hu, Fan Yang, Lintao Zhang, and Lidong Zhou. 2020. Rammer: Enabling Holistic Deep Learning Compiler Optimizations with rTasks. In USENIX OSDI. [33] MiniMax. 2025. MiniMax-M1: Scaling Test-Time Compute Efficiently with Lightning Attention. In arXiv. [34] Mooncake Team. 2025. Deploying Kimi K2 with PD Disaggregation and Large-Scale Expert Parallelism on 128 H200 GPUs. https://www. lmsys.org/blog/2025-07-20-k2-large-scale-ep/. (2025). [35] NVIDIA. 2025. CUDA Programming Guide. https://docs.nvidia.com/ cuda/cuda-c-programming-guide. (2025). [36] NVIDIA. 2026. Green Contexts. https://docs.nvidia.com/cuda/cudaprogramming-guide/04-special-topics/green-contexts.html. (2026). [37] NVIDIA Corporation. 2024. NVIDIA Multi-Instance GPU User Guide. https://docs.nvidia.com/datacenter/tesla/mig-user-guide/. (2024). [38] Pratyush Patel, Esha Choukse, Chaojie Zhang, Íñigo Goiri, Aashaka Shah, Saeed Maleki, and Ricardo Bianchini. 2023. Splitwise: Efficient generative LLM inference using phase splitting. In arXiv. [39] Ruoyu Qin, Zheming Li, Weiran He, Jialei Cui, Heyi Tang, Feng Ren, Teng Ma, Shangming Cai, Yineng Zhang, Mingxing Zhang, Yongwei
[1] 2025. CUDA Multi-Process Service. (2025). https://docs.nvidia.com/ deploy/pdf/CUDA_Multi_Process_Service_Overview.pdf. [2] Amey Agrawal, Nitin Kedia, Ashish Panwar, Jayashree Mohan, Nipun Kwatra, Bhargav Gulavani, Alexey Tumanov, and Ramachandran Ramjee. 2024. Taming Throughput-Latency tradeoff in LLM inference with Sarathi-Serve. In USENIX OSDI. [3] DeepSeek AI. 2026. Deepseek-v4: Towards highly efficient milliontoken context intelligence. arXiv (2026). [4] DeepSeek AI. 2026. Revealing the challenges of attention-ffn disaggregation for modern moe models and hardware systems. arXiv (2026). [5] Minimax AI. 2026. The minimax-m2 series: Mini activations unleashing max real-world intelligence. arXiv (2026). [6] Zhipu AI. 2026. Glm-5: from vibe coding to agentic engineering. arXiv (2026). [7] Zhihao Bai, Zhen Zhang, Yibo Zhu, and Xin Jin. 2020. Pipeswitch: Fast pipelined context switching for deep learning applications. In USENIX OSDI. [8] Tom B. Brown, Benjamin Mann, Nick Ryder, Melanie Subbiah, Jared Kaplan, Prafulla Dhariwal, Arvind Neelakantan, Pranav Shyam, Girish Sastry, Amanda Askell, Sandhini Agarwal, Ariel Herbert-Voss, Gretchen Krueger, Tom Henighan, Rewon Child, Aditya Ramesh, Daniel M. Ziegler, Jeffrey Wu, Clemens Winter, Christopher Hesse, Mark Chen, Eric Sigler, Mateusz Litwin, Scott Gray, Benjamin Chess, Jack Clark, Christopher Berner, Sam McCandlish, Alec Radford, Ilya Sutskever, and Dario Amodei. 2020. Language Models are Few-Shot Learners. (2020). [9] Li-Wen Chang, Wenlei Bao, Qi Hou, Chengquan Jiang, Ningxin Zheng, Yinmin Zhong, Xuanrun Zhang, Zuquan Song, Chengji Yao, Ziheng Jiang, Haibin Lin, Xin Jin, and Xin Liu. 2024. FLUX: Fast Softwarebased Communication Overlap On GPUs Through Kernel Fusion. In arXiv. [10] Yukang Chen, Weihao Cui, Han Zhao, Ziyi Xu, Xiaoze Fan, Xusheng Chen, Yangjie Zhou, Shixuan Sun, Bingsheng He, and Quan Chen. 2026. Towards High-Goodput LLM Serving with Prefill-decode Multiplexing. In ACM ASPLOS. [11] Xinhao Cheng, Zhihao Zhang, Yu Zhou, Jianan Ji, Jinchen Jiang, Zepeng Zhao, Ziruo Xiao, Zihao Ye, Yingyi Huang, Ruihang Lai, Hongyi Jin, Bohan Hou, Mengdi Wu, Yixin Dong, Anthony Yip, Zihao Ye, Songting Wang, Wenqin Yang, Xupeng Miao, Tianqi Chen, and Zhihao Jia. 2026. MPK: A Compiler and Runtime for Mega-Kernelizing Tensor Programs. In USENIX OSDI. [12] Patrick H. Coppock, Brian Zhang, Eliot H. Solomon, Vasilis Kypriotis, Leon Yang, Bikash Sharma, Dan Schatzberg, Todd C. Mowry, and Dimitrios Skarlatos. 2025. LithOS: An Operating System for Efficient Machine Learning on GPUs. In ACM SOSP. [13] DeepSeek-AI. 2024. DeepSeek-V3 Technical Report. In arXiv. [14] DeepSeek-AI. 2025. Day 6: One More Thing, DeepSeek-V3/R1 Inference System Overview. https://github.com/deepseek-ai/openinfra-index/blob/main/202502OpenSourceWeek/day_6_one_more_ thing_deepseekV3R1_inference_system_overview.md. (2025). [15] DeepSeek-AI. 2025. DeepEP: an efficient expert-parallel communication library. https://github.com/deepseek-ai/DeepEP. (2025). [16] Ruwen Fan, Tingxu Ren, Minhui Xie, Shiwei Gao, Jiwu Shu, and Youyou Lu. 2025. GPREEMPT: GPU Preemptive Scheduling Made General and Efficient. In USENIX ATC. [17] GLM-4.5 Team. 2025. GLM-4.5: Agentic, Reasoning, and Coding (ARC) Foundation Models. In arXiv. [18] Wentao Guo, Mayank Mishra, Xinle Cheng, Ion Stoica, and Tri Dao. 2025. SonicMoE: Accelerating MoE with IO and Tile-aware Optimizations. In arXiv. [19] Mingcong Han, Hanze Zhang, Rong Chen, and Haibo Chen. 2022. Microsecond-scale Preemption for Concurrent GPU-accelerated DNN Inferences. In USENIX OSDI. 13
Information Processing Systems. [58] Yinmin Zhong, Shengyu Liu, Junda Chen, Jianbo Hu, Yibo Zhu, Xuanzhe Liu, Xin Jin, and Hao Zhang. 2024. DistServe: Disaggregating Prefill and Decoding for Goodput-optimized Large Language Model Serving. In USENIX OSDI. [59] Kan Zhu, Yufei Gao, Yilong Zhao, Liangyu Zhao, Gefei Zuo, Yile Gu, Dedong Xie, Tian Tang, Qinyu Xu, Zihao Ye, Keisuke Kamahori, ChienYu Lin, Ziren Wang, Stephanie Wang, Arvind Krishnamurthy, and Baris Kasikci. 2025. NanoFlow: Towards Optimal Large Language Model Serving Throughput. In USENIX OSDI. [60] Ruidong Zhu, Ziheng Jiang, Chao Jin, Peng Wu, Cesar A. Stuardo, Dongyang Wang, Xinlei Zhang, Huaping Zhou, Haoran Wei, Yang Cheng, Jianzhe Xiao, Xinyi Zhang, Lingjun Liu, Haibin Lin, Li-Wen Chang, Jianxi Ye, Xiao Yu, Xuanzhe Liu, Xin Jin, and Xin Liu. 2025. MegaScale-Infer: Serving Mixture-of-Experts at Scale with Disaggregated Expert Parallelism. In arXiv.
Wu, Weimin Zheng, and Xinran Xu. 2025. Mooncake: A KVCachecentric Disaggregated Architecture for LLM Serving. ACM Trans. Storage (2025). [40] Qwen Team. 2025. Qwen3 Technical Report. In arXiv. [41] Xiaoxiang Shi, Colin Cai, Junjia Du, Zhanda Zhu, Xingda Wei, and Zhihao Jia. 2025. Nexus: Taming Throughput-Latency Tradeoff in LLM Serving via Efficient GPU Sharing. arXiv (2025). [42] Benjamin Spector, Jordan Juravsky, Stuart Sul, Owen Dugan, Dylan Lim, Dan Fu, Simran Arora, and Christopher Re. 2025. Look Ma, No Bubbles! Designing a Low-Latency Megakernel for Llama-1B. https: //hazyresearch.stanford.edu/blog/2025-05-27-no-bubbles. (2025). [43] StepFun. 2025. Step-3 is Large yet Affordable: Model-system Co-design for Cost-effective Decoding. In arXiv. [44] Xun Sun, Shaoyuan Chen, Pingchuan Ma, Yue Chen, Ziwei Yuan, Zhanhao Cao, Han Han, Shangming Cai, Teng Ma, Xuchun Shang, et al. 2026. Surviving Partial Rank Failures in Wide Expert-Parallel MoE Inference. arXiv (2026). [45] Tile-AI. 2025. TileRT: Tile-Based Runtime for Ultra-Low-Latency LLM Inference. https://github.com/tile-ai/TileRT. (2025). [46] Chao Wang, Pengfei Zuo, Zhangyu Chen, Yunkai Liang, Zhou Yu, and Ming-Chang Yang. 2025. Prefill-Decode Aggregation or Disaggregation? Unifying Both for Goodput-Optimized LLM Serving. arXiv (2025). [47] Xinming Wei, Chao Jin, Tuo Dai, Yinmin Zhong, Shan Yu, Chengxu Yang, Bingyang Wu, Zili Zhang, Jing Mai, Qianchao Zhu, et al. 2026. UltraEP: Unleash MoE Training and Inference on Rack-Scale Nodes with Near-Optimal Load Balancing. arXiv (2026). [48] Bingyang Wu, Shengyu Liu, Yinmin Zhong, Peng Sun, Xuanzhe Liu, and Xin Jin. 2024. Loongserve: Efficiently serving long-context large language models with elastic sequence parallelism. In ACM SOSP. [49] Bingyang Wu, Zili Zhang, Zhihao Bai, Xuanzhe Liu, and Xin Jin. 2023. Transparent GPU Sharing in Container Clouds for Deep Learning Workloads. In USENIX NSDI. [50] Bingyang Wu, Zili Zhang, Yinmin Zhong, Guanzhe Huang, Yibo Zhu, Xuanzhe Liu, and Xin Jin. 2025. TokenLake: A Unified Segment-level Prefix Cache Pool for Fine-grained Elastic Long-Context LLM Serving. arXiv (2025). [51] Xiaomi LLM-Core Team. 2026. MiMo-V2-Flash Technical Report. In arXiv. [52] Mingshu Zhai, Jiaao He, Zixuan Ma, Zan Zong, Runqing Zhang, and Jidong Zhai. 2023. SmartMoE: Efficiently Training Sparsely-Activated Models through Combining Offline and Online Parallelization. In USENIX ATC. [53] Shulai Zhang, Quan Chen, Weihao Cui, Han Zhao, Chunyu Xue, Zhen Zheng, Wei Lin, and Minyi Guo. 2025. Improving GPU Sharing Performance through Adaptive Bubbleless Spatial-Temporal Sharing. In EuroSys. [54] Shulai Zhang, Ningxin Zheng, Haibin Lin, Ziheng Jiang, Wenlei Bao, Chengquan Jiang, Qi Hou, Weihao Cui, Size Zheng, Li-Wen Chang, Quan Chen, and Xin Liu. 2025. Comet: Fine-grained Computationcommunication Overlapping for Mixture-of-Experts. In arXiv. [55] Zhexiang Zhang, Ye Wang, Yumiao Zhao, Jiayu Xiao, Qianjing Yang, Xiangyu Wang, Jingzhe Jiang, Qizhen Weng, Ruichuan Chen, Shaohuai Shi, Adel N. Toosi, Yin Chen, and Minchen Yu. 2026. Janus: Disaggregating Attention and Experts for Scalable MoE Inference. In arXiv. [56] Chenggang Zhao, Zhean Xu, Liang Zhao, Jiashi Li, Chenhao Xu, Anyi Xu, Shengyu Liu, Kexing Zhou, and Kuai Yu. 2025. DeepGEMM: clean and efficient BLAS kernel library on GPU. https: //github.com/deepseek-ai/DeepGEMM. (2025). [57] Lianmin Zheng, Liangsheng Yin, Zhiqiang Xie, Chuyue Sun, Jeff Huang, Cody Hao Yu, Shiyi Cao, Christos Kozyrakis, Ion Stoica, Joseph E. Gonzalez, Clark Barrett, and Ying Sheng. 2024. SGLang: Efficient Execution of Structured Language Model Programs. In Neural 14