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:
| Property | MoE Experts | Attention |
|---|---|---|
| Model Footprint | >95% of parameters | <5% of parameters |
| Computation Pattern | Dynamically sparse | Dense, structured |
| Parallelism | Expert Parallelism (EP) | Data Parallelism (DP) or Tensor Parallelism (TP) |
| Communication | All-to-all dispatch/combine | All-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:
- Bounded Preemption: Preemption bounded by one tile execution time (~25.3 μs) plus one cluster check epoch, independent of total operation length
- Spatial Multiplexing: Different CTA clusters execute different phases concurrently
- Temporal Multiplexing: Clusters switch phases at tile boundaries
- CUDA Graph Compatibility: No CPU intervention or kernel relaunch needed
The preemption mechanism uses a hierarchical decision propagation through memory hierarchy:
- System-scope flag signals urgent decode
- CTA 0 checks at tile boundaries, writes device-scope words
- Clusters read only their 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 and decode SM budgets to maximize goodput:
The tile-aware latency model fits:
where MoE footprint uses tile-aware token count:
Online SM reallocation adapts to runtime load:
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):
| Setting | ChunkedPrefill | Colocated | PDD | PDMux | ExpertPlex | Best Speedup |
|---|---|---|---|---|---|---|
| MiniMax/ShareGPT | 1.7 | 3.0 | 3.8 | 4.7 | 11.3 | 5.65× |
| MiniMax/LooGLE | - | 2.1 | - | 3.4 | 8.7 | 4.12× |
| GLM/ShareGPT | 0.6 | 1.0 | - | 1.5 | 1.5 | 3.30× |
| GLM/LooGLE | 0.3 | 0.6 | - | 0.9 | 1.5 | 5.00× |
Comparison to GPU Sharing Mechanisms:
| Mechanism | CUDA Graph | Temporal | Spatial | Bounded Preemption | Bounded Reallocation |
|---|---|---|---|---|---|
| API Interception | ✓ | ✓ | ✗ | ✗ | ✗ |
| CUDA Streams | ✓ | ✓ | ✗ | ✗ | ✗ |
| NVIDIA MPS | ✓ | ✗ | ✓ | ✗ | ✗ |
| Green Context | ✓ | ✗ | ✓ | ✗ | ✗ |
| NVIDIA MIG | ✓ | ✗ | Limited | ✗ | ✗ |
| 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
- Hybrid disaggregation-colocation architecture that matches phase demands with fewer GPUs and less traffic
- Adaptive Persistent Kernels providing tile-level preemption with sequence-length-independent bounds
- Attention-initiated one-sided communication eliminating deadlocks and enabling cross-phase overlap
- 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
- OasisKV: Scaling In-Decode KV Cache Beyond HBM with Lookahead Sparse Prefetching
OasisKV uses speculative decoding's draft tokens to predict future KV-cache access, enabling asynchronous prefetching that boosts LLM inference throughput up to 2.1x with negligible accuracy loss.
- Rethinking Self-Evolving Agents: Do We Still Need Prescribed Optimization Pipelines?
Frontier optimizers can compose task-specific improvement strategies online without prescribed pipelines, matching or beating them on 12 of 14 settings while using a third of the compute.
- RestoreKV: Recovering Full-Cache Behavior Under Aggressive Query-Agnostic KV Cache Eviction
RestoreKV generates a small, context-conditioned restore cache via LoRA-adapted tokens to complement evicted KV pairs, boosting RULER accuracy by up to 35 points at 5% budget with negligible overhead.