MoonEP

MoonEP is an open-source expert-parallel communication library for mixture-of-experts models, released by Moonshot AI on July 27, 2026 under the MIT License.^[1][3][4] Its stated design goal is unusual for the category: rather than making an imbalanced all-to-all faster, MoonEP removes the imbalance, guaranteeing that every expert-parallel rank receives exactly S x K tokens (S input tokens per rank, K routed experts per token) no matter how skewed the router output is, by planning a small number of duplicated "dynamic redundant experts" on the GPU at every step.^[1] The library was published during Moonshot's Kimi K3 Open Day alongside the Kimi K3 weights and technical report, and Section 5.2.1 of that report describes MoonEP as the expert-parallel scheme used for K3's 2.8-trillion-parameter pre-training.^[2][17][18]

Scroll sideways for more →

Background: why expert parallelism needs its own communication library

In a sparse MoE transformer, the feed-forward sub-layer is replaced by many expert networks and a router that sends each token to its top-K experts. Under expert parallelism, those experts are split across devices, so computing one MoE layer requires two collectives: a dispatch all-to-all that ships each token to the ranks holding its chosen experts, and a combine all-to-all that returns the expert outputs to the token's home rank. That is a different axis from data parallelism, which replicates the model and splits the batch, tensor parallelism, which splits weight matrices, and pipeline parallelism, which splits layers across stages. Production runs combine all of them: K3 pre-training used PP with virtual stages, EP, ZeRO-1 data parallelism, pipeline ZeRO-2 gradient sharding, and context parallelism together.^[2][10]

The awkward part is that routing is data-dependent. Some experts are far more popular than others, so the rank hosting a hot expert receives more tokens, computes longer, and makes every other rank wait at the next synchronization point, and that straggler sets the step time for the whole job. Dynamic per-rank token counts cause a second problem: activation shapes change every step and layer, forcing a host-device synchronization to learn the shapes before launching expert kernels and fragmenting GPU memory, since buffers must be provisioned for the worst case.^[2][12]

Plainly stated: expert parallelism is a mail-sorting problem. Each step the router addresses millions of tokens and the network delivers them, and if one sorting office gets triple the mail, everyone else waits. Existing libraries make the trucks faster. MoonEP instead prints a temporary copy of the overloaded office and hands it to whoever is idle.

Measuring imbalance: maxvio

The metric MoonEP sweeps in its benchmarks is maxvio (maximal violation), introduced by DeepSeek researchers in the 2024 paper on auxiliary-loss-free load balancing as (max_i Load_i - mean Load) / mean Load, where Load_i is the number of tokens assigned to expert i and the mean is the expected load under perfect balance.^[14] MoonEP's README writes the same quantity as max_e (T_e / T_bar) - 1.^[1] A maxvio of 0 means perfect balance; a maxvio of 10 means the hottest expert carries 11 times the balanced load. The harness reverse-solves a lognormal routing distribution to hit target maxvio values of 0.2, 1, 10, and 20.^[8]

How MoonEP works

Dynamic redundant experts and the E/R bound

Replicating hot experts is not new; what MoonEP adds is a guarantee. The Kimi K3 technical report proves that for any router output, a plan exists in which every rank receives exactly S x K tokens using at most E/R redundant experts per rank, where E is the number of routed experts and R is the EP group size (Theorem 1), and that the bound is essentially tight, since there are router outputs requiring ceil(E(R-1)/R^2) redundant experts, which approaches E/R for large R (Theorem 2).^[2] The proof is constructive: repeatedly fill an underloaded rank from an overloaded one until it holds exactly S x K tokens, which terminates after at most R-1 fills and leaves each rank drawing its remote tokens from a single other rank, hence from at most E/R distinct remote experts.^[2]

That bound is the engineering payoff. Because reserving E/R slots per rank always admits a feasible plan, training can never be interrupted by an unplannable step. The report contrasts this with ECHO and UltraEP, which preset the number of redundant experts or impose a per-rank token cap, so training stops when no feasible plan fits the cap, and the cap needs manual tuning while still leaving residual imbalance.^[2]

Weight and gradient buffers

MoonEP's contract with a training framework is one contiguous symmetric-memory weight tensor per expert projection (gate, up, down) shaped [E+B, H, H'], laid out identically on every rank, plus a planner-produced cu_seqlens array that tells the grouped GEMM which expert rows are active this step.^[1] Rows [0, E) are all ranks' local experts, E/R rows each, and each chunk physically is the home rank's parameter memory mapped everywhere through symmetric memory. Rows [E, E+B) are the local prefetch slots that prefetch_weight fills, backed by a process-global pool shared across layers, so the extra memory cost is B expert weights per projection in total rather than per layer.^[1]

Training must use B = E/R so every expert the grouped GEMM touches is local; inference needs no gradients and may use fewer slots, with B = 3 to 4 recommended and any overflow read directly from the home rank through the symmetric mapping, slightly slower but still correct.^[1] Training mirrors the weight layout with an fp32 gradient buffer, except that prefetch-slot gradients are backed by a separate reduce buffer so duplicated experts' gradients stay invisible to the framework's own gradient reduction; reduce_grad then has each rank read its own experts' slots from every rank's reduce buffer over NVLink, accumulate them into its local parameter gradient, and zero the consumed slots.^[1]

Online planning

Deciding which experts to duplicate, and where, happens every layer of every micro-batch, so it has to be nearly free. Moonshot computed exact solutions offline with integer linear programming for representative cases as a reference, then built a GPU planning kernel it describes as near-optimal with negligible overhead that always respects the E/R bound.^[2] The implementation is a single cooperative grid launch running four phases in one kernel, written in the CUTLASS Python DSL, with a software grid barrier across blocks and a system-scope self-resetting atomic barrier across ranks on the NVLink metadata buffer.^[6] Dispatch is likewise CuTe DSL, using warp specialization and cp.async.bulk TMA transfers; the C++ extension is limited to virtual-memory allocation, IPC handle exchange, and multicast setup.^[7][11]

Zero copy, static shapes, and sync-free execution

Because the planner precomputes each token's destination, the permute that groups tokens by expert is fused into the communication itself: tokens are written straight into their expert-grouped positions on remote ranks, and views of the communication buffer are handed to the expert FFN, so there is no permute-in, no permute-out, and no copy from the communication buffer to a user buffer.^[1][2] With zero_copy=True on both dispatch and combine, the FFN reads and writes the buffer in place; the views alias state that the next call overwrites, so they must not be held across communication calls or saved for autograd.^[1]

Two consequences follow. First, buffer size: the report notes that the same copy-free path in DeepEP would need a communication buffer of S x K x R under worst-case imbalance, while perfect balance lets MoonEP allocate a fixed S x K.^[2] Second, shapes are statically known at every layer, which removes both the per-layer host-device synchronization that dropless MoE implementations need to learn their computation shapes and the fragmentation that comes with ever-changing activation shapes.^[1][2]

API and requirements

The public surface is a single Buffer object with four collective operations, each of which accepts async_finish=True to run on the communication stream and return a CUDA event.^[1]

Scroll sideways for more →

Installation is pip install -e ., which compiles the moonep._C extension; the six test suites (planning, dispatch, combine, end-to-end, gradient reduce, prefetch) run under torchrun --nproc_per_node=8 and require multiple GPUs plus NVLink.^[1][5] The code asserts that num_ep_ranks equals the process-group world size and that E is divisible by R; num_sms defaults to 32 and B to E // num_ep_ranks.^[9] A contributor reported that three kernels also require H and H' to be multiples of 128, which the README's weight contract does not state.^[13]

Reported performance

The README publishes no numeric benchmark tables. Both comparisons against DeepEP v2 are presented as plots, one for communication time and one for end-to-end training iteration time, both run on H20 GPUs at EP=8 while sweeping router imbalance.^[1] The claims, all Moonshot's own, are that MoonEP's communication time is consistently below DeepEP v2 at every imbalance level and stays almost flat as maxvio grows, while DeepEP v2, whose latency is set by the hottest rank, degrades steadily; that the bars already include MoonEP's extra planning and weight-prefetch kernels, leaving total dispatch on par with DeepEP v2's dispatch alone and combine significantly faster at every level; and that end-to-end iteration time stays flat with no memory fragmentation, whereas DeepEP's climbs and training eventually runs out of memory at high imbalance.^[1]

The communication benchmark ships as benchmarks/bench_vs_deepep.py, which asserts a world size of 8 and defaults to 8,192 tokens per rank, 384 experts, hidden size 7,168, top-8 routing, an expert intermediate dimension of 2,048, 32 SMs, token padding of 128, 20 warmup iterations, and 50 timed iterations, with both libraries fed identical routing matrices and tensors under one timing harness. Those defaults match Kimi K2's MoE shape (384 routed experts, hidden 7,168, top-8, expert intermediate 2,048) rather than K3's.^[2][8] The script compares only DeepEP v2's elastic path with expanded dispatch, since DeepEP v1 groups received tokens by source rank and performs the expert permute outside the library.^[8] No end-to-end training script is included, and the README does not state the token count behind the end-to-end figure. A user raised that as issue #6 on July 28, 2026, noting that prefetching weights over NVLink costs more than reading them from local HBM, so the point where balance repays the prefetch depends on the tokens-per-rank regime; the issue was still unanswered on July 31, 2026.^[13]

Relationship to Kimi K3

Kimi K3 is a 2.78-trillion-parameter MoE model with 104.2 billion parameters active per token, 896 routed experts, 16 experts selected per token, and two shared experts.^[2] Fine-grained routing at that ratio makes the dispatch and combine collectives, and their imbalance, a first-order cost, which is the setting MoonEP was built for: Moonshot's open-day summary, relayed by Chinese outlets, describes it as a high-performance communication library built for very large fine-grained MoE that keeps expert-parallel communication efficient even when expert load is badly skewed.^[17] Moonshot's K3 tech blog from the July 16 announcement promised the weights by July 27 but named no infrastructure libraries.^[23] The technical report lists MoonEP among K3's infrastructure contributions and states that it provides "perfectly balanced expert execution with static computation shapes and zero-copy communication" for the 2.8T-parameter pre-training run.^[2]

One widely repeated framing needs care. The report attributes its approximately 2.5-fold improvement in overall scaling efficiency over Kimi K2 to Kimi Delta Attention, attention residuals, Stable LatentMoE, and refined data and training recipes, and lists MoonEP separately under infrastructure; the figure is not presented as a MoonEP result.^[2] The report also does not disclose K3's EP degree, its cluster size, or a measured end-to-end gain attributable to MoonEP alone.^[2]

Comparison with DeepEP, ECHO, and UltraEP

Scroll sideways for more →

The README's Echo credit is a link to arXiv:2603.07685, NVIDIA's Megatron Core MoE paper, where ECHO is described as one technique rather than a standalone system.^[1][12] The AcclEP acknowledgment was added to the README on 2026-07-28, the day after the announcement.^[1] The "Zhenwu PPU" listed as under review refers to Alibaba's T-Head accelerator line, which the company describes as a parallel processing unit designed for AI training and inference.^[21]

Sibling releases at Kimi K3 Open Day

Chinese coverage of the open day described three infrastructure technologies published with the K3 weights, covering high-performance communication, high-performance operators, and distributed reinforcement-learning environments.^[17][18] Two need clarification. FlashKDA, the CUTLASS-based kernel library for Kimi Delta Attention, was not new: its repository has been public since 2026-04-20 with a design write-up dated 2026-04-22, and IT Home noted it had been open-sourced earlier.^[18][20] AgentENV, the platform for running agent environments at scale behind K3's agentic RL training, comes from kvcache-ai, the organization behind Mooncake, not the MoonshotAI org; its repository was created on 2026-07-23 under the MIT license.^[19]

Reception, limitations, and open questions

Coverage in the week after release stayed in the developer press: MarkTechPost summarized the library on July 29, 2026, and a July 30 analysis argued that enterprise adopters will want the side-by-side throughput curves and cost-per-token figures the release does not provide.^[15][24] By July 31, 2026 the repository had 949 stars, 99 forks, and 22 open issues, none with a maintainer reply.^[1][13]

Several limitations are visible in the code and the issue tracker rather than in the announcement:

• Scale-up only, so far. Everything shipped assumes NVLink symmetric memory within a node, the benchmark asserts a world size of 8, and a user asked on July 30 whether the library is NVLink-only and what to use for EP across nodes.^[8][13]
• Framework integration is invasive. MoonEP requires one contiguous [E+B, H, H'] symmetric-memory tensor per expert projection per layer, with contiguity a hard requirement, so it is not a drop-in replacement for an existing dispatcher.^[1]
• Weight prefetch must be hidden. A reader traced the communication-stream ordering and asked whether the design assumes two-batch overlap at the framework level to hide prefetch_weight behind the expert FFN; the question is unanswered.^[13]
• Early-stage packaging. The package is version 0.0.1, pins an exact CuTe DSL version, has no CI, and uses assertions rather than exceptions for user-facing validation, all filed as issues within a day of release.^[5][13]
• No independent benchmarks. As of July 31, 2026 the comparison against DeepEP v2 rests on Moonshot's own plots and reproduction script; no third-party evaluation had been published, and neither vLLM nor SGLang had announced an integration.^[1][15][24]

MoonEP's significance is that it turns MoE load balancing from a best-effort heuristic into a provable invariant, then spends that invariant on things balance alone does not buy: static shapes, no host synchronization, no fragmentation, and a communication buffer that does not grow with EP size. Whether the trade, extra NVLink traffic for duplicated expert weights in exchange for perfect balance, pays off outside Moonshot's own distributed training regime is what the release leaves open.

See also

• Expert parallelism
• DeepEP
• Mixture of experts
• Kimi K3
• Moonshot AI
• Distributed training
• NVLink
• Megatron-LM

References

1. ^Moonshot AI, "MoonEP: A Perfectly Balanced Expert Parallelism Library via Dynamic Redundant Experts", GitHub repository and README, initial commit 2026-07-24, AcclEP acknowledgment added 2026-07-28. github.com/...MoonEP. Accessed 2026-07-31.
2. ^Moonshot AI, "Kimi K3: Open Frontier Intelligence", technical report (PDF), Section 5.2.1 and Appendix E, 2026-07-27. github.com/...k3_tech_report.pdf. Accessed 2026-07-31.
3. ^Kimi.ai (@Kimi_Moonshot), post announcing MoonEP, X, 2026-07-27. x.com/...2081763086281973847. Accessed 2026-07-31.
4. ^Moonshot AI, MoonEP LICENSE (MIT License, Copyright 2026 Moonshot AI), GitHub. github.com/...LICENSE. Accessed 2026-07-31.
5. ^Moonshot AI, MoonEP `setup.py` (version 0.0.1, `nvidia-cutlass-dsl==4.4.2`), GitHub. github.com/...setup.py. Accessed 2026-07-31.
6. ^Moonshot AI, MoonEP `moonep/planning.py` (CuTe DSL planning kernel), GitHub. github.com/...planning.py. Accessed 2026-07-31.
7. ^Moonshot AI, MoonEP `moonep/dispatch.py` (warp-specialized TMA dispatch), GitHub. github.com/...dispatch.py. Accessed 2026-07-31.
8. ^Moonshot AI, MoonEP `benchmarks/bench_vs_deepep.py`, GitHub. github.com/...bench_vs_deepep.py. Accessed 2026-07-31.
9. ^Moonshot AI, MoonEP `moonep/api.py` (Buffer construction and constraints), GitHub. github.com/...api.py. Accessed 2026-07-31.
10. ^DeepSeek-AI, "DeepEP: an efficient expert-parallel communication library", GitHub repository and README, created 2025-02-17. github.com/...DeepEP. Accessed 2026-07-31.
11. ^Moonshot AI, MoonEP `moonep/buffer.py` (NVLink symmetric-memory allocation and IPC handle exchange), GitHub. github.com/...buffer.py. Accessed 2026-07-31.
12. ^Zijie Yan et al. (NVIDIA), "Scalable Training of Mixture-of-Experts Models with Megatron Core", arXiv:2603.07685, submitted 2026-03-08, revised 2026-03-10 (Section 4.3.7, "ECHO: Elastic Cloning for Hot Experts"). arxiv.org/...2603.07685. Accessed 2026-07-31.
13. ^MoonEP issue tracker, issues #1 to #22 (including #6 on end-to-end benchmark token counts, #15 on hiding prefetch latency, #19 on the H and H' multiple-of-128 rule, and #22 on NVLink-only operation), GitHub, 2026-07-28 to 2026-07-30. github.com/...issues. Accessed 2026-07-31.
14. ^Lean Wang, Huazuo Gao, Chenggang Zhao, Xu Sun, Damai Dai, "Auxiliary-Loss-Free Load Balancing Strategy for Mixture-of-Experts", arXiv:2408.15664, 2024-08-28 (defines MaxVio in Section 4.1). arxiv.org/...2408.15664. Accessed 2026-07-31.
15. ^Michal Sutter, "Moonshot AI Open-Sources MoonEP: A Perfectly Balanced Expert Parallelism Library for MoE Training", MarkTechPost, 2026-07-29. marktechpost.com/...elism-library-for-moe-training Accessed 2026-07-31.
16. ^Dots-Infra, "UltraEP: Production-ready MoE load balancing via real-time expert replication", GitHub repository and README, v1.0.0 released 2026-07-15. github.com/...UltraEP. Accessed 2026-07-31.
17. ^Sina Finance, "月之暗面开源Kimi K3模型权重、技术报告，公开三大关键Infra技术", 2026-07-28. finance.sina.com.cn/...doc-inikimee8072440.shtml. Accessed 2026-07-31.
18. ^IT Home (IT之家), "月之暗面开源 Kimi K3 模型，2.8 万亿参数", 2026-07-27. ithome.com/...259.htm. Accessed 2026-07-31.
19. ^kvcache-ai, "AgentENV (AENV) is a distributed platform for running agent environments at scale", GitHub repository, created 2026-07-23, MIT license. github.com/...AgentENV. Accessed 2026-07-31.
20. ^Moonshot AI, "FlashKDA: Flash Kimi Delta Attention, high-performance KDA kernels built on CUTLASS", GitHub repository and README, created 2026-04-20. github.com/...FlashKDA. Accessed 2026-07-31.
21. ^South China Morning Post, "Alibaba's T-Head unit unveils details of AI chip designed to rival Nvidia's GPUs", 2026-01-29. scmp.com/...s-ai-chip-designed-rival-nvidias-gpus. Accessed 2026-07-31.
22. ^DeepSeek-AI, pull request #605, "[Public release 26/04] Introducing EPv2: faster EP, and Engram/PP/CP supports", merged 2026-04-29. github.com/...605. Accessed 2026-07-31.
23. ^Moonshot AI, "Kimi K3" tech blog, kimi.com, July 2026. kimi.com/...kimi-k3. Accessed 2026-07-31.
24. ^i10x, "Moonshot AI Open-Sources MoonEP for MoE Training", 2026-07-30. i10x.ai/...t-ai-moonep-moe-networking-bottlenecks. Accessed 2026-07-31.