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

Summary (Overview)

  • Hybrid Architecture: ExpertPlex introduces a novel hybrid disaggregation-colocation architecture that shares massive MoE expert weights across prefill and decode phases while disaggregating lightweight attention modules, eliminating over 95% of duplicate model weights.
  • Adaptive Persistent Kernels (APKs): A new GPU scheduling mechanism that provides tile-level preemption and reallocation with bounds independent of sequence length, supporting both spatial and temporal multiplexing without CPU intervention.
  • Attention-Initiated One-Sided Communication: Replaces conventional two-sided MoE communication with one-sided push/pull primitives that eliminate deadlock risks, reduce network interference, and enable cross-phase communication-computation overlap.
  • Cross-Stack Placement Optimizer: Jointly optimizes placement, parallelism, overlap strategies, and APK scheduling policies to maximize goodput, achieving up to 2.01× improvement over instance-level PDD and 1.66× over Green Context-based colocation.
  • Significant Performance Gains: Evaluated on MiniMax-M2.7 and GLM-5.1-FP8, ExpertPlex improves goodput by up to 5.65× over chunked prefill and 4.12× over naive colocated serving.

Introduction and Theoretical Foundation

Background

Modern Large Language Models (LLMs) increasingly adopt sparse Mixture-of-Experts (MoE) architectures that dynamically activate a subset of experts per token. While this scales model capacity without proportional computation increase, it creates significant serving challenges. The paper notes that expert weights now dominate model footprints: 95% in DeepSeek-V4-Pro, 96% in GLM-5.1-FP8, and 98% in MiniMax-M2.7.

Problem with Existing Approaches

Instance-Level Prefill-Decode Disaggregation (PDD) places phases on separate full-model replicas:

  • Coarse allocation granularity (e.g., one DeepSeek-V3 deployment uses 32 prefill + 320 decode GPUs per unit)
  • Duplicated expert weights displace KV-cache capacity
  • Large failure blast radius with hierarchical communication
  • Inflexible scaling in 100+ GPU increments

Prefill-Decode Colocation (via Green Context) partitions each GPU by phase but:

  • Cannot adapt to dynamic MoE load variations across layers
  • Suffers from head-of-line blocking or resource bubbles
  • Forces wider parallelism, increasing communication overhead
  • Leaves cross-phase network interference unmanaged

Theoretical Foundation

The key insight is that MoE weights and attention have fundamentally different resource characteristics:

PropertyMoE ExpertsAttention
Model Footprint>95% of parameters<5% of parameters
Computation PatternDynamically sparseDense, structured
ParallelismExpert Parallelism (EP)Data Parallelism (DP) or Tensor Parallelism (TP)
CommunicationAll-to-all dispatch/combineAll-reduce or none

This asymmetry suggests a hybrid approach: share the massive expert weights while independently scaling attention modules per phase.

Methodology

System Architecture

ExpertPlex organizes GPUs into three types of servers:

  • Prefill attention servers: Handle attention computation for prefill
  • Decode attention servers: Handle attention computation for decode
  • MoE servers: Execute expert computation for both phases

Adaptive Persistent Kernels (APKs)

APKs schedule MoE computation at tile granularity (the smallest independently completable unit of grouped GEMMs), providing:

  1. Bounded Preemption: Preemption bounded by one tile execution time (~25.3 μs) plus one cluster check epoch, independent of total operation length
  2. Spatial Multiplexing: Different CTA clusters execute different phases concurrently
  3. Temporal Multiplexing: Clusters switch phases at tile boundaries
  4. CUDA Graph Compatibility: No CPU intervention or kernel relaunch needed

The preemption mechanism uses a hierarchical decision propagation through memory hierarchy:

  • System-scope flag PP signals urgent decode
  • CTA 0 checks PP at tile boundaries, writes device-scope words i:Pi\forall i: \mathcal{P}_i
  • Clusters read only their Pi\mathcal{P}_i via DSMEM broadcast
  • mbarrier handoff coordinates warps within CTAs

Attention-Initiated One-Sided Communication

Replaces two-sided ring-buffer protocols with one-sided primitives:

Dispatch (Push): Attention servers write activations directly to final MoE buffers via NVLink peer stores or one-sided RDMA writes, publishing a ready signal only after payload visibility.

Combine (Pull): A single-thread WaitDone kernel on the attention GPU observes completion, then pulls results via NVLink loads or RDMA reads.

Traffic Isolation: Prefill scale-out traffic routes through prefill attention servers (hierarchical path) while decode communicates directly, with InfiniBand virtual-lane priorities as a fallback.

Cross-Stack Placement Optimizer

The optimizer jointly searches over layouts L\mathcal{L} and decode SM budgets QQ_\ell to maximize goodput:

G(,q)=min(BpTp,BdTdOˉ)(1)G(\ell, q) = \min\left(\frac{B_p}{T_p}, \frac{B_d}{T_d \bar{O}}\right) \tag{1}

The tile-aware latency model fits:

t^c(x,s)=αc+βcx+γcxs+δcxs2(2)\hat{t}_c(x, s) = \alpha_c + \beta_c x + \gamma_c x s + \delta_c x s^2 \tag{2}

where MoE footprint uses tile-aware token count:

xmoe=eme>0meMt(3)x_{\mathrm{moe}} = \sum_{e | m_e > 0} \left\lceil \frac{m_e}{M_t} \right\rceil \tag{3}

Online SM reallocation adapts to runtime load:

q=min(Qmax,qxmoexmoec)(4)q' = \min\left(Q_{\max}, \left\lceil \frac{q \cdot x_{\mathrm{moe}}}{x_{\mathrm{moe}}^\star} \right\rceil_c\right) \tag{4}

Empirical Validation / Results

Experimental Setup

  • Models: MiniMax-M2.7 (230GB FP8, single-node) and GLM-5.1-FP8 (756GB FP8, multi-node)
  • Hardware: NVIDIA H800 GPUs with NVLink and 200 Gbps InfiniBand
  • Workloads: ShareGPT (short) and LooGLE (long) input/output distributions
  • Baselines: SGLang-ChunkedPrefill, SGLang-Colocated, SGLang-PDD, SGLang-PDMux

Key Results

End-to-End Goodput (requests/sec/node):

SettingChunkedPrefillColocatedPDDPDMuxExpertPlexBest Speedup
MiniMax/ShareGPT1.73.03.84.711.35.65×
MiniMax/LooGLE-2.1-3.48.74.12×
GLM/ShareGPT0.61.0-1.51.53.30×
GLM/LooGLE0.30.6-0.91.55.00×

Comparison to GPU Sharing Mechanisms:

MechanismCUDA GraphTemporalSpatialBounded PreemptionBounded Reallocation
API Interception
CUDA Streams
NVIDIA MPS
Green Context
NVIDIA MIGLimited
ExpertPlex

Overhead Analysis:

  • APK scheduling: <12% overhead for contiguous layout, <20 μs for masked layout
  • Communication: within 5% of DeepEP v1 in normal mode
  • Preemption intervals: all <25.3 μs for GEMMs, <10.7 μs for data processing

Theoretical and Practical Implications

Architectural Insights

The paper demonstrates that the optimal disaggregation boundary is at the MoE-attention interface, not at the model level. This insight enables fine-grained resource allocation in single-GPU units while avoiding massive weight duplication. The work provides a principled framework for reasoning about where to draw disaggregation boundaries based on parameter fraction versus computational characteristics.

GPU Resource Management

APKs establish that tile-level scheduling can provide production-grade isolation and preemption while maintaining near-native performance. The bound on preemption independent of sequence length is a significant theoretical contribution, enabling predictable latency guarantees for latency-sensitive workloads sharing GPUs with throughput-oriented work.

Communication Protocol Design

The attention-initiated one-sided communication model shows that removing receiver-side coordination can improve both correctness (deadlock freedom) and performance. The hierarchical traffic isolation strategy offers a template for managing heterogeneous QoS requirements over shared networks.

Practical Deployment Implications

  • Small clusters: Can realize target prefill:decode ratios previously requiring hundreds of GPUs
  • Fault containment: Smaller deployment units reduce failure blast radius
  • Elastic scaling: Fine-grained (single-GPU) scaling matches traffic shifts efficiently
  • Memory efficiency: Eliminates >95% duplicate weights, freeing memory for KV cache

Conclusion

ExpertPlex presents a compelling case for rethinking disaggregation boundaries in LLM serving systems. By sharing MoE experts across phases while disaggregating attention, it achieves the isolation benefits of PDD without its resource waste, and the memory efficiency of colocation without its interference problems.

Key Contributions

  1. Hybrid disaggregation-colocation architecture that matches phase demands with fewer GPUs and less traffic
  2. Adaptive Persistent Kernels providing tile-level preemption with sequence-length-independent bounds
  3. Attention-initiated one-sided communication eliminating deadlocks and enabling cross-phase overlap
  4. Cross-stack placement optimizer jointly optimizing placement, parallelism, and scheduling

Future Directions

  • Extending APK scheduling to other sparse computation patterns beyond MoE
  • Exploring adaptive tile sizes that trade off scheduling granularity for kernel efficiency
  • Investigating the interaction between the proposed techniques and multi-tenant/multi-tenant serving
  • Applying the hybrid disaggregation principle to other model architectures with heterogeneous weight distributions

The work establishes that fine-grained, workload-aware resource management can deliver order-of-magnitude improvements in serving efficiency, providing a foundation for future systems that must serve increasingly large and complex model architectures.

Related papers