DeepEP README

The project's README from github.com/deepseek-ai/DeepEP. As one markdown file: README.md.

DeepEP

DeepEP (DeepEveryParallel) is a high-performance communication library for modern machine learning training and inference. The library currently focuses on expert parallelism (EP) — providing high-throughput and low-latency all-to-all GPU kernels (MoE dispatch and combine) with low-precision support including FP8 — while also offering experimental primitives for pipeline parallelism (PP), context parallelism (CP), and remote memory access (Engram), all designed for zero or minimal SM occupation. All kernels are compiled at runtime via a lightweight Just-In-Time (JIT) module, requiring no CUDA compilation during installation.

Despite its lightweight design, DeepEP's performance matches or exceeds hardware bandwidth limits across various configurations.

News

New features

Notes

Still on-going features

For the legacy V1 documentation (NVSHMEM-based), see docs/legacy.md.

Performance

Following V3's configuration, we tested with 8K tokens per batch, 7168 hidden dimensions, top 8 experts, FP8 dispatching, and BF16 combining, and obtained the following results:

Arch NIC type Topo Dispatch Bottleneck Bandwidth Combine Bottleneck Bandwidth #SMs
SM90 CX7 EP 8 x 2 90 GB/s (RDMA) 81 GB/s (RDMA) 12
SM90 CX7 EP 8 x 4 61 GB/s (RDMA) 61 GB/s (RDMA) 6
SM100 CX7 EP 8 x 2 90 GB/s (RDMA) 91 GB/s (RDMA) 12
SM100 N/A EP 8 726 GB/s (NVLink) 740 GB/s (NVLink) 64 (Max perf)
SM100 N/A EP 8 643 GB/s (NVLink) 675 GB/s (NVLink) 24 (Min #SM)

Notes: the results are logical bandwidth. For example, under the EP 8 x 2 case, 90 GB/s actually contains local rank traffic.

Comparing with V1, V2 achieves up to 1.3x peak performance, while saving up to 4x SM count.

We omit results for larger EP configurations for the time being, but encourage interested users to benchmark them directly. Based on our internal experience, we expect the kernel to continue saturating hardware bandwidth at scale.

For V1 performance data, see docs/legacy.md.

Quick start

Requirements

Install NCCL dependency

We recommend using pip to install NCCL so that DeepEP can automatically locate it within the Python environment. You can install it using the following command:

pip install "nvidia-nccl-cu13>=2.30.4" --no-deps

Install NVSHMEM dependency

DeepEP also depends on NVSHMEM to provide support for legacy methods. Please refer to our NVSHMEM Installation Guide for instructions.

Development

# Build and make symbolic links for SO files
python setup.py build
# You may modify the specific SO names according to your own platform
ln -s build/lib.linux-x86_64-cpython-38/deep_ep_cpp.cpython-38-x86_64-linux-gnu.so

# Run test cases
# NOTES: you may modify the `init_dist` function in `tests/utils/envs.py`
# according to your own cluster settings, and launch into multiple nodes
python tests/elastic/test_ep.py
python tests/elastic/test_agrs.py
python tests/elastic/test_engram.py
python tests/elastic/test_pp.py

Installation

python setup.py install

Then, import deep_ep in your Python project, and enjoy!

Interfaces and examples

Buffer initialization

In V2, all EP operations — high-throughput and low-latency — are unified under a single ElasticBuffer interface. The buffer can be initialized by specifying MoE settings directly, and the optimal SM and QP counts are calculated analytically.

import torch
import torch.distributed as dist
from typing import Optional

from deep_ep import ElasticBuffer

# Communication buffer (will allocate at runtime)
_buffer: Optional[ElasticBuffer] = None

# Number of SMs to use for communication kernels (will be set at buffer creation)
_num_comm_sms: int = 0


def get_buffer(group: dist.ProcessGroup,
               num_max_tokens_per_rank: int,
               hidden: int,
               num_topk: int,
               num_experts: int,
               use_fp8_dispatch: bool = False) -> ElasticBuffer:
    """Initialize or retrieve the ElasticBuffer for EP communication."""
    global _buffer, _num_comm_sms

    # Check if we can reuse the existing buffer
    required_bytes = ElasticBuffer.get_buffer_size_hint(
        group, num_max_tokens_per_rank, hidden,
        num_topk=num_topk, use_fp8_dispatch=use_fp8_dispatch,
    )
    if _buffer is not None and _buffer.group == group and _buffer.num_bytes >= required_bytes:
        return _buffer

    # Allocate a new buffer with MoE settings
    # NOTES: V2 buffer size consumption is larger than V1
    _buffer = ElasticBuffer(
        group,
        num_max_tokens_per_rank=num_max_tokens_per_rank,
        hidden=hidden,
        num_topk=num_topk,
        use_fp8_dispatch=use_fp8_dispatch,
    )

    # V2 analytically calculates the optimal SM count — no more auto-tuning needed
    # You may also specify `num_sms` manually in dispatch/combine calls to override
    _num_comm_sms = _buffer.get_theoretical_num_sms(num_experts, num_topk)

    return _buffer

Example use in model training or inference prefilling

V2 unifies the dispatch and combine APIs into a single ElasticBuffer interface. The example below shows how to use them for training (with backward passes) or inference prefilling.

import torch
import torch.distributed as dist
from typing import Tuple, Union

from deep_ep import ElasticBuffer, EPHandle, EventOverlap


def dispatch_forward(x: Union[torch.Tensor, Tuple[torch.Tensor, torch.Tensor]],
                     topk_idx: torch.Tensor, topk_weights: torch.Tensor,
                     num_experts: int,
                     num_max_tokens_per_rank: int,
                     expert_alignment: int = 1) -> \
        Tuple[Union[torch.Tensor, Tuple[torch.Tensor, torch.Tensor]],
              torch.Tensor, torch.Tensor, EPHandle, EventOverlap]:
    """
    MoE dispatch: route tokens to the corresponding experts across all ranks.
    Supports both BF16 and FP8 (x as a tuple of [data, scale_factors]) inputs.
    """
    global _buffer, _num_comm_sms

    recv_x, recv_topk_idx, recv_topk_weights, handle, event = _buffer.dispatch(
        x,
        topk_idx=topk_idx,
        topk_weights=topk_weights,
        num_experts=num_experts,
        num_max_tokens_per_rank=num_max_tokens_per_rank,
        expert_alignment=expert_alignment,
        num_sms=_num_comm_sms,
        async_with_compute_stream=True,
    )

    # `handle` contains routing metadata for the subsequent combine call
    # `handle.num_recv_tokens_per_expert_list` provides per-expert token counts for GEMM
    # Use `event.current_stream_wait()` to synchronize the compute stream before using results
    return recv_x, recv_topk_idx, recv_topk_weights, handle, event


def dispatch_backward(grad_recv_x: torch.Tensor,
                      grad_recv_topk_weights: torch.Tensor,
                      handle: EPHandle) -> Tuple[torch.Tensor, torch.Tensor, EventOverlap]:
    """The backward pass of MoE dispatch is actually a combine."""
    global _buffer, _num_comm_sms

    combined_grad_x, combined_grad_topk_weights, event = _buffer.combine(
        grad_recv_x,
        handle=handle,
        topk_weights=grad_recv_topk_weights,
        num_sms=_num_comm_sms,
        async_with_compute_stream=True,
    )

    return combined_grad_x, combined_grad_topk_weights, event


def combine_forward(x: torch.Tensor,
                    handle: EPHandle) -> Tuple[torch.Tensor, EventOverlap]:
    """MoE combine: reduce expert outputs back to their original ranks."""
    global _buffer, _num_comm_sms

    combined_x, _, event = _buffer.combine(
        x,
        handle=handle,
        num_sms=_num_comm_sms,
        async_with_compute_stream=True,
    )

    return combined_x, event


def combine_backward(grad_combined_x: Union[torch.Tensor, Tuple[torch.Tensor, torch.Tensor]],
                     handle: EPHandle) -> \
        Tuple[Union[torch.Tensor, Tuple[torch.Tensor, torch.Tensor]], EventOverlap]:
    """The backward pass of MoE combine is actually a dispatch."""
    global _buffer, _num_comm_sms

    grad_x, _, _, _, event = _buffer.dispatch(
        grad_combined_x,
        handle=handle,
        num_sms=_num_comm_sms,
        async_with_compute_stream=True,
    )

    return grad_x, event

For communication-computation overlap, use the EventOverlap interface to manage dependencies between the communication stream and the compute stream:

# After dispatch, overlap computation while communication is in-flight
recv_x, recv_topk_idx, recv_topk_weights, handle, event = dispatch_forward(...)

# ... do some independent computation here ...

# Wait for communication to finish before using results
event.current_stream_wait()

# Now safe to use recv_x, recv_topk_idx, recv_topk_weights

Example use in inference decoding

For inference decoding, the same ElasticBuffer is used. The handle-caching pattern allows reusing routing metadata across iterations when the gating decisions remain unchanged, avoiding redundant CPU synchronization.

import torch
from typing import Tuple, Optional, Union

from deep_ep import ElasticBuffer, EPHandle, EventOverlap


def decode_dispatch(x: Union[torch.Tensor, Tuple[torch.Tensor, torch.Tensor]],
                    topk_idx: torch.Tensor, topk_weights: torch.Tensor,
                    num_experts: int,
                    num_max_tokens_per_rank: int,
                    cached_handle: Optional[EPHandle] = None) -> \
        Tuple[Union[torch.Tensor, Tuple[torch.Tensor, torch.Tensor]],
              torch.Tensor, torch.Tensor, EPHandle, EventOverlap]:
    """
    MoE dispatch for inference decoding.
    If `cached_handle` is provided, the layout is reused without CPU synchronization.
    """
    global _buffer, _num_comm_sms

    if cached_handle is not None:
        # Reuse cached handle: skip layout recomputation and CPU sync
        recv_x, _, _, handle, event = _buffer.dispatch(
            x,
            handle=cached_handle,
            num_sms=_num_comm_sms,
            async_with_compute_stream=True,
        )
        return recv_x, cached_handle.topk_idx, None, handle, event

    recv_x, recv_topk_idx, recv_topk_weights, handle, event = _buffer.dispatch(
        x,
        topk_idx=topk_idx,
        topk_weights=topk_weights,
        num_experts=num_experts,
        num_max_tokens_per_rank=num_max_tokens_per_rank,
        num_sms=_num_comm_sms,
        async_with_compute_stream=True,
    )

    return recv_x, recv_topk_idx, recv_topk_weights, handle, event


def decode_combine(x: torch.Tensor,
                   handle: EPHandle) -> Tuple[torch.Tensor, EventOverlap]:
    """MoE combine for inference decoding."""
    global _buffer, _num_comm_sms

    combined_x, _, event = _buffer.combine(
        x,
        handle=handle,
        num_sms=_num_comm_sms,
        async_with_compute_stream=True,
    )

    return combined_x, event

Environment variables

The library provides some environment variables, which may be useful:

Some environment variables are persistent: they are captured at build time and baked into the installed package as default values. At import time, these defaults are applied automatically unless overridden by current environment variables. The persistent variables are: EP_JIT_CACHE_DIR, EP_JIT_PRINT_COMPILER_COMMAND, EP_NUM_TOPK_IDX_BITS, EP_NCCL_ROOT_DIR.

For additional details, please refer to the test code or review the corresponding Python documentation.

Network configurations

DeepEP is fully tested with InfiniBand networks. However, it is theoretically compatible with RDMA over Converged Ethernet (RoCE) as well.

Traffic isolation

Traffic isolation is supported by InfiniBand through Virtual Lanes (VL).

To prevent interference between different types of traffic, we recommend segregating workloads across different virtual lanes as follows:

For DeepEP V2, you can control the virtual lane assignment by setting the sl_idx argument or the EP_OVERRIDE_RDMA_SL environment variable.

Adaptive routing

Adaptive routing is an advanced routing feature provided by InfiniBand switches that can evenly distribute traffic across multiple paths. Even though adaptive routing introduces additional latency, we still recommend enabling it under all network load conditions.

Congestion control

Congestion control is disabled because it hurts maximum bandwidth. If congestion is unavoidable in some scenarios, we recommend assigning those workloads to low-priority virtual lanes.

PCI atomic mode

If the hardware supports it, we recommend using the following command to set the NIC's PCI_ATOMIC_MODE to improve RDMA atomic operation performance:

sudo mlxconfig -y -d mlx5_$i set PCI_ATOMIC_MODE=4

Experimental branches

Community forks

Acknowledgement

DeepEP V2 is built on top of the NCCL Gin backend. Thanks to @sjeaugey, @pakmarkthub, @sb17v, @xiaofanl-nvidia, and the NCCL team for their support!

License

This code repository is released under the MIT License.

Citation

@misc{deepep2025,
      title={DeepEP: an efficient expert-parallel communication library},
      author={Chenggang Zhao and Shangyan Zhou and Liyue Zhang and Chengqi Deng and Zhean Xu and Yuxuan Liu and Kuai Yu and Jiashi Li and Liang Zhao},
      year={2025},
      publisher = {GitHub},
      howpublished = {\url{https://github.com/deepseek-ai/DeepEP}},
}

DeepEP (deepseek-ai/DeepEP)

Searching Database…

10.3K ·GitHub
DeepEP
Что это такое?DeepEP — это высокопроизводительная библиотека для коммуникаций в моделях MoE (mixture-of-experts). Она предоставляет быстрые GPU-ядра all-to-all для операций dispatch и combine в параллелизме экспертов, которые используются при обучении и инференсе моделей машинного обучения.
Это фреймворк для обучения или библиотека, которую я подключаю к существующей настройке PyTorch?Это библиотека, а не полноценный фреймворк: вы импортируете её в код обучения или инференса на PyTorch, чтобы ускорить коммуникации в параллелизме экспертов. Ядра компилируются во время выполнения через лёгкий JIT-модуль, а не заранее, как при классической компиляции CUDA (ahead-of-time).
Как установить и начать пользоваться?Установите зависимость NCCL через pip, чтобы DeepEP мог сам её найти, затем просто выполните import deepep в своем Python-проекте. Отдельного шага сборки CUDA нет, потому что ядра компилируются с JIT при первом использовании.
Какие GPU и версии ПО нужны?Нужны GPU Hopper (SM90) или другое оборудование, совместимое с SM90-PTX, Python 3.8+, CUDA 12.3+, PyTorch 2.10+ и NCCL 2.30.4+, а также NVLink для внутриузловой связи и RDMA-сеть для межузловой.
Работает ли он на GPU, отличных от Hopper, и поддерживаются ли в нем все функции V1?Официально он рассчитан на GPU Hopper (SM90); более широкая поддержка (например, AMD) есть только в форках от сообщества. Кроме того, в V2 не поддерживается режим low-latency RDMA EP с 0 SM из V1, а Engram, PP и CP помечены как экспериментальные.
Что изменилось между V1 и V2?V2 — это полная переработка, которая использует в несколько раз меньше потоковых мультипроцессоров (SM), переходит с бэкенда NVSHMEM на NCCL Gin и объединяет API для высокой пропускной способности и низкой задержки в единый интерфейс ElasticBuffer, который поддерживает масштабирование до EP2048.
Что такое бэкенд NCCL Gin?Это новый, более легковесный бэкенд, который DeepEP V2 использует вместо NVSHMEM. Он также позволяет DeepEP переиспользовать существующие коммуникаторы NCCL приложения, вместо того чтобы создавать отдельные.
Что такое функции Engram, PP и CP?Это экспериментальные примитивы, которые работают без использования SM и построены вместе с параллелизмом экспертов: Engram отвечает за удалённый доступ к памяти (через RDMA), PP — за конвейерный параллелизм (через RDMA), а CP — за параллелизм контекста (через Copy Engine).
Могу ли я запускать DeepEP на GPU не от NVIDIA?Официально нет, но форки от сообщества расширяют его возможности: ROCm/mori переносит бэкенд на GPU AMD, а uccl/uccl-ep позволяет запускать DeepEP на разнородных GPU и сетевых адаптерах (NIC).
What are similar projects?DeepGEMM, DeepEP-Ascend, TensorRT-LLM, DeepGEMM-Ascend, taskflow, tilelang, DeepSpeed, ColossalAI
Which topics is it in?hpc, distributed-computing, header-only, jit, gpu-programming, mixture-of-experts, moe, high-performance, gpu-acceleration, distributed-systems, inference, nvidia
Creator
deepseek-ai
Size
3.8 MB
Stars
10.3K
Updated
CommitOct 15Releasev1.2.1Sep 16
Age
1y
Commits · 90 Days