ExpertPlex: A High-Goodput Disaggregated Serving System for MoE LLMs with Adaptive Persistent Kernels

Abstract: 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\% 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) attention-initiated 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 MiniMax-M2.7 and GLM-5.1-FP8 show that ExpertPlex improves goodput by up to 2.01$\times$ over instance-level prefill-decode disaggregation and 1.66$\times$ over prefill-decode colocation.

• The paper introduces ExpertPlex, a hybrid serving architecture that disaggregates attention while sharing MoE experts, avoiding costly expert-weight replication and enabling finer-grained scaling.
• Adaptive Persistent Kernels schedule prefill and decode work at tile granularity, limiting decode overhead to 8% while delivering near-exclusive decode latency and 1.12× prefill slowdown in kernel tests.
• The system improves P90 goodput by up to 5.65× over chunked prefill, 2.01× over instance-level PDD, and 4.12× over colocated serving across MiniMax-M2.7 and GLM-5.1-FP8 workloads.

Motivation and problem statement

ExpertPlex addresses the serving inefficiency that arises when Mixture-of-Experts (MoE) LLMs are served under conventional prefill-decode disaggregation (PDD) or prefill-decode colocation. The authors observe that expert weights now dominate model footprints — 95% in DeepSeek-V4-Pro, 96% in GLM-5.1-FP8, and 98% in MiniMax-M2.7 — so any scheme that replicates the full model per phase (instance-level PDD) duplicates a massive memory footprint. Reported deployments make this concrete: a DeepSeek-V3 unit combines 32 prefill and 320 decode GPUs, another uses 176 GPUs, and Kimi-K2 has been deployed on 128 H200s. Such coarse deployment units prevent small clusters from realizing the target prefill-to-decode ratio, force large clusters to scale in hundreds-of-GPU increments, and widen the failure blast radius because hierarchical all-to-all communication couples many ranks into one group.

Colocation via Green Context spatial partitioning avoids duplication but cannot follow the temporal dynamics of MoE serving: expert load varies across ranks and layers, attention and MoE modules have different demands within a layer, and each MoE module alternates between dispatch, computation, and combine. Fixed partitions therefore cause either head-of-line blocking — a prefill grouped GEMM can run tens to hundreds of milliseconds while a decode GEMM finishes in hundreds of microseconds, an 84–101× duration gap measured on MiniMax-M2.7 with EP4 — or idle reserved SMs (resource bubbles). Partitioning every GPU also reduces each phase's local resources, forcing wider parallelism, more communication, and unmanaged cross-phase network interference.

Architecture: hybrid disaggregation-colocation

The central design is a hybrid boundary: share MoE experts across phases, but disaggregate attention modules. Since attention holds under 5% of parameters, disaggregating attention onto whole GPUs per phase avoids duplicating expert weights while giving each phase full local compute capacity, reducing the required degree of parallelism and communication relative to intra-GPU partitioning. Each node assigns subsets of GPUs to prefill attention servers, decode attention servers, and shared MoE servers; prefill or decode server sets may be empty. This MoE-weight-independent allocation boundary enables finer elastic scaling and smaller failure blast radius than instance-level PDD.

Adaptive Persistent Kernel (APK)

Sharing MoE GPUs requires finer control than existing mechanisms provide. ExpertPlex runs an Adaptive Persistent Kernel per MoE GPU that schedules grouped GEMMs and data processing at tile granularity — the smallest independently completable unit, with boundaries occurring every 2.2–25.3 μs regardless of operation length. Because phases switch only at committed tile boundaries, no accumulator, TMA transaction, or communication state remains live, so preemption requires no checkpoint, restore, or recomputation. The kernel preserves native DeepGEMM-style warp specialization, TMA multicast, and CUDA Graph compatibility, and all scheduling occurs on the GPU without CPU intervention.

Bounded preemption is nontrivial because high-performance kernels pipeline warps and CTAs across tiles; naive per-warp switching would deadlock on mbarrier or TMA-multicast dependencies. APK propagates a single cooperative decision down the memory hierarchy — a system-scope word read by CTA 0, per-cluster device-scope words, DSMEM broadcast by the cluster leader, and an mbarrier handoff across warps — so the slowest CTA finishes at most its current tile before the cluster converges. The preemption bound is one tile execution plus one local check epoch, independent of sequence length.

A microbenchmark against CUDA stream priorities, MPS, and Green Context on concurrent prefill- and decode-shaped grouped GEMMs (GLM-5.1-FP8, 8192 vs. 128 tokens) substantiates the claim: under a decode-latency protection target, priority streams inflate decode latency 13.79× over exclusive execution, while MPS and Green Context keep decode near-exclusive but slow prefill 3.33× and 4.07× respectively; APK adds only 8% decode overhead and slows prefill by 1.12×. APK is thus the only evaluated mechanism occupying the low-decode-latency, high-prefill-throughput corner of the Pareto frontier. Tile-level scheduling itself costs under 12% on prefill contiguous layouts and under 20 μs on decode masked layouts. The authors compare APK's preemption interval (all MoE operations below 25.3 μs, GEMMs below 10.7 μs) with prior preemption systems (REEF's best reported 35 μs with recomputation), while conceding those systems were not designed for MoE workloads and lack TMA multicast, CTA cluster, and CUDA Graph support — so the comparison is a reference point rather than a controlled evaluation.

Attention-initiated one-sided MoE communication

Conventional two-sided all-to-all requires receiver-side progress; with phases independently scheduled, some MoE ranks may run prefill kernels while others run decode, and each phase waits for receiver work on ranks occupied by the other — a cross-phase deadlock cycle. Reserving polling SMs avoids this only by wasting compute. ExpertPlex instead removes MoE-side coordination entirely: APK preallocates final dispatch and combine buffers (a consequence of persistent execution requiring fixed templates and descriptors), and attention servers push activations via NVLink peer stores or one-sided RDMA writes, then publish ready signals; combine is a pull initiated from the attention side through a single-thread WaitDone kernel that coexists with TBO/SBO overlap. Because prefill and decode follow independent dependency chains, APK can overlap one phase's communication with the other's computation — an opportunity unavailable to intra-phase overlap schemes.

Traffic isolation exploits topology: prefill scale-out traffic is routed between prefill attention servers (one RDMA transfer per node, then NVLink multicast, preserving hierarchical deduplication), while decode communicates directly attention-to-MoE. When a destination node lacks a prefill server, direct RDMA is used with prefill assigned a lower-priority InfiniBand virtual lane. Measured against DeepEP v1 on 16 GPUs, the one-sided path tracks normal-mode dispatch/combine within about 5% and stays within roughly 45 μs of DeepEP in low-latency mode, so removing MoE-side coordination does not sacrifice communication efficiency.

Cross-stack placement optimizer

Because placement, parallelism, overlap strategy, and the decode SM budget interact, ExpertPlex jointly searches them offline. The latency model fits a small per-component form $$\hat{t}_c(x,s)=\alpha_c+\beta_c x+\gamma_c xs+\delta_c xs^2$$ from profiling on a few GPUs, with a key observation that MoE latency tracks executed tiles rather than raw token count: the MoE footprint sums per-expert ceilings $\lceil m_e/M_t\rceil$, since each active expert triggers at least one tile regardless of token volume. The optimizer maximizes request-level goodput $G(\ell,q)=\min(B_p/T_p,\; B_d/(T_d\bar{O}))$ over SLO-feasible configurations. Online, APK treats $q$ as a contention policy, scaling the decode SM budget with the ratio of current to expected MoE footprint, giving decode priority and prefill the remainder. The model's accuracy depends on the fitted quadratic form and profiled-GPU representativeness; the paper does not report prediction error separately from end-to-end results.

Evaluation

ExpertPlex is implemented by modifying DeepGEMM, DeepEP v1, and SGLang, evaluated on MiniMax-M2.7 (single-node H800, 230 GB FP8) and GLM-5.1-FP8 (up to three nodes, 756 GB FP8), with ShareGPT and LooGLE workloads under Poisson arrivals, using P90 goodput (TTFT + TPOT SLOs) as the primary metric. Headline results:

Workload	vs. ChunkedPrefill	vs. Colocated	vs. PDD	vs. PDMux
MiniMax-M2.7 / ShareGPT	5.65×	2.72×	2.01×	1.41×
MiniMax-M2.7 / LooGLE	— (fails SLO)	4.12×	—	1.28×
GLM-5.1-FP8 / ShareGPT	3.3×	1.5×	not run (OOM)	~parity
GLM-5.1-FP8 / LooGLE	5.0×	2.5×	not run (OOM)	1.66×

Two results deserve emphasis. First, GLM-5.1-FP8 could not run under PDD at all within 24 GPUs due to memory duplication of expert weights — a direct demonstration of the memory-efficiency claim. Second, ExpertPlex and SGLang-PDMux achieve roughly equal goodput (~1.5 req/s/node) on GLM-5.1-FP8 with ShareGPT: PDMux's tensor-parallel attention gives short requests a TTFT advantage, but this fades on LooGLE, where long requests expose network interference and Green Context's utilization limits. The advantage of ExpertPlex is therefore workload-dependent, concentrated on long-input and high-contention regimes.

Limitations and open questions

Several constraints qualify the results. The evaluation uses H800 hardware with NVLink and InfiniBand; the one-sided communication design depends on NVLink peer access and RDMA one-sided primitives (IBGDA), and its portability to other interconnects or to GPUs with different cluster/DSMEM capabilities is not established. The offline optimizer targets an expected input/output length distribution, so performance under sharp traffic shifts relies on online SM reallocation rather than re-placement, and the paper does not quantify adaptation under distribution drift. The preemption comparison against LithOS, REEF, GPreempt, and related systems is acknowledged as non-controlled, since those systems lack support for modern kernel features. Finally, the evaluation covers two models with 256 routed experts and top-8 activation; whether the tile-footprint latency model and the hybrid boundary remain optimal for models with different expert granularity, shared-expert ratios, or attention architectures (only full attention and DSA are tested) is left open.

Conclusion

ExpertPlex combines a hybrid disaggregation-colocation boundary, adaptive persistent kernels with sequence-length-independent bounded tile-level preemption, attention-initiated one-sided MoE communication, and a cross-stack placement optimizer. On MiniMax-M2.7 and GLM-5.1-FP8 it improves P90 goodput by up to 5.65× over chunked prefill, 2.01× over instance-level PDD, 1.66× over Green Context-based colocation, and 4.12× over naive colocation, while adding at most 8% decode overhead and 1.12× prefill slowdown in isolated kernel benchmarks. The system demonstrates that phase isolation and weight sharing need not be mutually exclusive when the disaggregation boundary is placed at the attention–MoE interface and GPU sharing is pushed to tile granularity.