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

> ExpertPlex shares MoE expert weights across prefill and decode while disaggregating attention, achieving up to 5.65x goodput gains via tile-level preemptive GPU scheduling.

- **Source:** [arXiv](https://arxiv.org/abs/2607.18002)
- **Published:** 2026-08-22
- **Permalink:** https://picx.dev/p/Un91AC
- **Whiteboard:** https://picx.dev/p/Un91AC/image

## Summary

# 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:

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 $P$ signals urgent decode
- CTA 0 checks $P$ at tile boundaries, writes device-scope words $\forall i: \mathcal{P}_i$
- Clusters read only their $\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 $\mathcal{L}$ and decode SM budgets $Q_\ell$ to maximize goodput:

$$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:

$$\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:

$$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\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)**:

| 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

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.

---

_Markdown view of https://picx.dev/p/Un91AC, served by PicX — AI-generated visual whiteboard summaries of research papers._