DeepEP Documentation

File: README.md Copy Chapter

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 Copy Chapter

• V2 release: A complete refactoring of Expert Parallelism — achieving extreme performance with several times fewer SM resources compared to V1, while supporting significantly larger scale-up and scale-out domains. V2 has also switched from the NVSHMEM backend to the more lightweight NCCL Gin backend.

New features Copy Section

• Fully JIT (Just-In-Time compilation)
• NCCL Gin backend
  • Header-only & lightweight
  • Able to reuse existing NCCL communicators
• EPv2
  • High-throughput and low-latency APIs unified into a single ElasticBuffer interface, with a new GEMM layout
  • Larger scale-up & scale-out domain support (up to EP2048)
  • Analytical SM & QP count calculation — no more auto-tuning needed
  • Both hybrid & direct modes remain supported
  • For V3-like legacy training, SM usage reduced from 24 to 4 - 6 while maintaining equivalent or better performance
• 0 SM Engram (with RDMA)
• 0 SM PP (with RDMA)
• 0 SM CP (with Copy Engine)

Notes Copy Section

• Buffer size consumption is larger than V1
• 0 SM RDMA low-latency EP is no longer supported
• Engram, PP, and CP are experimental features

Still on-going features Copy Section

• Elastic GPU & CPU buffers: A contiguous virtual address space that maps to a hybrid of GPU and CPU physical memory under the hood, enabling fully automatic and transparent Engram or imbalanced EP
• Reducing intermediate buffer sizes by leveraging EP replay to handle load imbalance
• All-gather updates and reduce-scatter implementations for DP & TP

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

Performance Copy Chapter

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:

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 Copy Chapter

Requirements Copy Section

• Hopper (SM90) GPUs, or other architectures with SM90 PTX ISA support
• Python 3.8 and above
• CUDA version
  • CUDA 12.3 and above for SM90 GPUs
• PyTorch 2.10 and above
• NCCL 2.30.4 and above
• NVLink for intranode communication
• RDMA network for internode communication

Install NCCL dependency Copy Section

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 Copy Section

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

Development Copy Section

# 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 Copy Section

python setup.py install

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

Interfaces and examples Copy Chapter

Buffer initialization Copy Section

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 Copy Section

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.

/* Detailed source-code truncated for AI context efficiency. */

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 Copy Section

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.

/* Detailed source-code truncated for AI context efficiency. */

Environment variables Copy Section

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

• General
  • EP_BUFFER_DEBUG: 0 or 1, print buffer initialization, SM approximation, and backend debugging information, 0 by default
  • EP_SUPPRESS_NCCL_CHECK: 0 or 1, suppress NCCL version mismatch checking, 0 by default
  • EP_AVOID_RECORD_STREAM: 0 or 1, avoid record_stream on output tensors, 0 by default
  • EP_NUM_TOPK_IDX_BITS: integer, override the number of bits for top-k index encoding, 0 (auto) by default
• Networking
  • EP_NIC_NAME: string, the default NIC name used to query NIC properties, mlx5_0 by default
  • EP_OVERRIDE_RDMA_SL: integer, override the RDMA service level index for traffic isolation
  • EP_DISABLE_GIN: 0 or 1, disable the NCCL Gin backend (fall back to non-Gin path), 0 by default
• JIT
  • EP_JIT_DEBUG: 0 or 1, print JIT debugging information, 0 by default
  • EP_JIT_CACHE_DIR: string, cache directory for compiled kernels, $HOME/.deep_ep by default
  • EP_JIT_NVCC_COMPILER: string, NVCC compiler path; defaults to torch.utils.cpp_extension.CUDA_HOME
  • EP_JIT_CPP_STANDARD: integer, C++ standard version, 20 by default
  • EP_JIT_PRINT_COMPILER_COMMAND: 0 or 1, print compilation commands, 0 by default
  • EP_JIT_PTXAS_VERBOSE: 0 or 1, show detailed PTXAS output, 0 by default
  • EP_JIT_PTXAS_CHECK: 0 or 1, assert no local memory usage in compiled kernels, 0 by default
  • EP_JIT_WITH_LINEINFO: 0 or 1, embed source line info for profiling tools, 0 by default
  • EP_JIT_DUMP_ASM: 0 or 1, dump both PTX and SASS, 0 by default
  • EP_JIT_DUMP_PTX: 0 or 1, dump PTX output, 0 by default
  • EP_JIT_DUMP_SASS: 0 or 1, dump SASS output, 0 by default
• Debug and profiling
  • EP_GIN_GDAKI_DEBUG: 0 or 1, enable NCCL Gin GDAKI debugging output, 0 by default
  • EP_USE_NVIDIA_TOOLS: 0 or 1, skip internal profiling when running under external NVIDIA tools, 0 by default
  • EP_DISABLE_BARRIER_PROFILING: 0 or 1, disable barrier-based communication profiling in benchmarks, 0 by default
• Build
  • EP_NCCL_ROOT_DIR: string, path to the NCCL installation directory; auto-detected from the Python environment if not set
  • EP_NVSHMEM_ROOT_DIR: string, path to the NVSHMEM installation directory; auto-detected from the Python environment if not set
  • TORCH_CUDA_ARCH_LIST: string, list of target CUDA architectures, e.g. "9.0"
  • DISABLE_SM90_FEATURES: 0 or 1, disable SM90 features for legacy methods, 0 by default
  • DISABLE_AGGRESSIVE_PTX_INSTRS: 0 or 1, disable aggressive load/store instructions in legacy methods, 0 by default

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 Copy Chapter

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

Traffic isolation Copy Section

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:

• expert-parallel workloads
• other workloads

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 Copy Section

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 Copy Section

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 Copy Section

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 Copy Chapter

• Zero-copy
  • Removing the copy between PyTorch tensors and communication buffers, which reduces the SM usages significantly for normal kernels
  • This PR is authored by Tencent Network Platform Department
• Eager
  • Using a low-latency protocol removes the extra RTT latency introduced by RDMA atomic OPs
• Hybrid-EP
  • A new backend implementation using TMA instructions for minimal SM usage and larger NVLink domain support
  • Fine-grained communication-computation overlap for single-batch scenarios
  • PCIe kernel support for non-NVLink environments
  • NVFP4 data type support
• AntGroup-Opt
  • This optimization series is authored by AntGroup Network Platform Department
  • Normal-SMFree Eliminating SM from RDMA path by decoupling comm-kernel execution from NIC token transfer, freeing SMs for compute
  • LL-SBO Overlapping Down GEMM computation with Combine Send communication via signaling mechanism to reduce end-to-end latency
  • LL-Layered Optimizing cross-node LL operator communication using rail-optimized forwarding and data merging to reduce latency
• Mori-EP
  • ROCm/AMD GPU support powered by MORI backend (low-latency mode)
• nvDev
  • V2-based branch with the latest CUDA features, such as Compute Fabric Transport (CFT) that brings better latency on small token sizes.

Community forks Copy Chapter

• uccl/uccl-ep - Enables running DeepEP on heterogeneous GPUs (e.g., Nvidia, AMD) and NICs (e.g., EFA, Broadcom, CX7)
• Infrawaves/DeepEP_ibrc_dual-ports_multiQP - Adds multi-QP solution and dual-port NIC support in IBRC transport
• antgroup/DeepXTrace - A diagnostic analyzer for efficient and precise localization of slow ranks
• ROCm/mori - AMD's next-generation communication library for performance-critical AI workloads (e.g., Wide EP, KVCache transfer, Collectives)

Acknowledgement Copy Chapter

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 Copy Chapter

This code repository is released under the MIT License.

Citation Copy Chapter

@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}},
}

File: docs/legacy.md Copy Chapter

DeepEP V1 (Legacy)

Note: This is the archived documentation for DeepEP V1 (NVSHMEM-based). For the latest V2 documentation, see the main README.

DeepEP (DeepEveryParallel) V1 is the original high-performance communication library for modern machine learning, focused on expert parallelism (EP). It provides high-throughput and low-latency all-to-all GPU kernels, which are also known as MoE dispatch and combine. The library also supports low-precision operations, including FP8.

To align with the group-limited gating algorithm proposed in the DeepSeek-V3 paper, DeepEP V1 offers a set of kernels optimized for asymmetric-domain bandwidth forwarding, such as forwarding data from NVLink domain to RDMA domain. These kernels deliver high throughput, making them suitable for both training and inference prefilling tasks. Additionally, they support SM (Streaming Multiprocessors) number control.

For latency-sensitive inference decoding, DeepEP V1 includes a set of low-latency kernels with pure RDMA to minimize delays. The library also introduces a hook-based communication-computation overlapping method that does not occupy any SM resource.

Notice: the implementation in this library may have some slight differences from the DeepSeek-V3 paper.

Performance Copy Chapter

Normal kernels with NVLink and RDMA forwarding Copy Section

We test normal kernels on H800 (~160 GB/s NVLink maximum bandwidth), with each connected to a CX7 InfiniBand 400 Gb/s RDMA network card (~50 GB/s maximum bandwidth). And we follow the DeepSeek-V3/R1 pretraining setting (4096 tokens per batch, 7168 hidden, top-4 groups, top-8 experts, FP8 dispatching and BF16 combining).

Low-latency kernels with pure RDMA Copy Section

We test low-latency kernels on H800 with each connected to a CX7 InfiniBand 400 Gb/s RDMA network card (~50 GB/s maximum bandwidth). And we follow a typical DeepSeek-V3/R1 production setting (128 tokens per batch, 7168 hidden, top-8 experts, FP8 dispatching and BF16 combining).

Quick Start Copy Chapter

Requirements Copy Section

• Ampere (SM80), Hopper (SM90) GPUs, or other architectures with SM90 PTX ISA support
• Python 3.8 and above
• CUDA version
  • CUDA 11.0 and above for SM80 GPUs
  • CUDA 12.3 and above for SM90 GPUs
• PyTorch 2.1 and above
• NVLink for intranode communication
• RDMA network for internode communication

Download and install NVSHMEM dependency Copy Section

DeepEP V1 depends on NVSHMEM. Please refer to the NVSHMEM Installation Guide for instructions.

Development Copy Section

# Build and make symbolic links for SO files
NVSHMEM_DIR=/path/to/installed/nvshmem 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.py`
# according to your own cluster settings, and launch into multiple nodes
python tests/test_intranode.py
python tests/test_internode.py
python tests/test_low_latency.py

Installation Copy Section

NVSHMEM_DIR=/path/to/installed/nvshmem python setup.py install

Installation environment variables

• NVSHMEM_DIR: the path to the NVSHMEM directory, disable all internode and low-latency features if not specified
• DISABLE_SM90_FEATURES: 0 or 1, whether to disable SM90 features, it is required for SM90 devices or CUDA 11
• TORCH_CUDA_ARCH_LIST: the list of target architectures, e.g. TORCH_CUDA_ARCH_LIST="9.0"
• DISABLE_AGGRESSIVE_PTX_INSTRS: 0 or 1, whether to disable aggressive load/store instructions, see Undefined-behavior PTX usage for more details

Network Configurations Copy Chapter

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

Traffic isolation Copy Section

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:

• workloads using normal kernels
• workloads using low-latency kernels
• other workloads

For DeepEP V1, you can control the virtual lane assignment by setting the NVSHMEM_IB_SL environment variable.

Adaptive routing Copy Section

Adaptive routing is an advanced routing feature provided by InfiniBand switches that can evenly distribute traffic across multiple paths. Enabling adaptive routing can completely eliminate network congestion caused by routing conflicts, but it also introduces additional latency. We recommend the following configuration for optimal performance:

• enable adaptive routing in environments with heavy network loads
• use static routing in environments with light network loads

Congestion control Copy Section

Congestion control is disabled as we have not observed significant congestion in our production environment.

Interfaces and Examples Copy Chapter

Example use in model training or inference prefilling Copy Section

The normal kernels can be used in model training or the inference prefilling phase (without the backward part) as the below example code shows.

/* Detailed source-code truncated for AI context efficiency. */

Moreover, inside the dispatch function, we may not know how many tokens to receive for the current rank. So an implicit CPU wait for GPU received count signal will be involved, as the following figure shows.

Example use in inference decoding Copy Section

The low latency kernels can be used in the inference decoding phase as the below example code shows.

/* Detailed source-code truncated for AI context efficiency. */

For two-micro-batch overlapping, you can refer to the following figure. With our receiving hook interface, the RDMA network traffic is happening in the background, without costing any GPU SMs from the computation part. But notice, the overlapped parts can be adjusted, i.e., the 4 parts of attention/dispatch/MoE/combine may not have the exact same execution time. You may adjust the stage settings according to your workload.

Roadmap (V1) Copy Chapter

• AR support
• Refactor low-latency mode AR code
• A100 support (intranode only)
• Support BF16 for the low-latency dispatch kernel
• Support NVLink protocol for intranode low-latency kernels
• TMA copy instead of LD/ST
  • Intranode kernels
  • Internode kernels
  • Low-latency kernels
• SM-free kernels and refactors
• Fully remove undefined-behavior PTX instructions

Notices Copy Chapter

Easier potential overall design

The V1 implementation uses queues for communication buffers which save memory but introduce complexity and potential deadlocks. If you're implementing your own version based on DeepEP V1, consider using fixed-size buffers allocated to maximum capacity for simplicity and better performance. For a detailed discussion of this alternative approach, see https://github.com/deepseek-ai/DeepEP/issues/39.

Undefined-behavior PTX usage

• For extreme performance, we discover and use an undefined-behavior PTX usage: using read-only PTX ld.global.nc.L1::no_allocate.L2::256B to read volatile data. The PTX modifier .nc indicates that a non-coherent cache is used. But the correctness is tested to be guaranteed with .L1::no_allocate on Hopper architectures, and performance will be much better. The reason we guess may be: the non-coherent cache is unified with L1, and the L1 modifier is not just a hint but a strong option, so that the correctness can be guaranteed by no dirty data in L1.
• Initially, because NVCC could not automatically unroll volatile read PTX, we tried using __ldg (i.e., ld.nc). Even compared to manually unrolled volatile reads, it was significantly faster (likely due to additional compiler optimizations). However, the results could be incorrect or dirty. After consulting the PTX documentation, we discovered that L1 and non-coherent cache are unified on Hopper architectures. We speculated that .L1::no_allocate might resolve the issue, leading to this discovery.
• If you find kernels not working on some other platforms, you may add DISABLE_AGGRESSIVE_PTX_INSTRS=1 to setup.py and disable this, or file an issue.

Auto-tuning on your cluster

For better performance on your cluster, we recommend to run all the tests and use the best auto-tuned configuration. The default configurations are optimized on the DeepSeek's internal cluster.

File: docs/nvshmem.md Copy Chapter

Install NVSHMEM

Important notices Copy Chapter

This project is neither sponsored nor supported by NVIDIA.

Use of NVIDIA NVSHMEM is governed by the terms at NVSHMEM Software License Agreement.

Prerequisites Copy Chapter

Hardware requirements:

• GPUs inside one node needs to be connected by NVLink
• GPUs across different nodes needs to be connected by RDMA devices, see GPUDirect RDMA Documentation
• InfiniBand GPUDirect Async (IBGDA) support, see IBGDA Overview
• For more detailed requirements, see NVSHMEM Hardware Specifications

Software requirements:

• NVSHMEM v3.3.9 or later

Installation procedure Copy Chapter

1. Install NVSHMEM binaries Copy Section

NVSHMEM 3.3.9 binaries are available in several formats:

• Tarballs for x86_64 and aarch64
• RPM and deb packages: instructions can be found on the NVSHMEM installer page
• Conda packages through conda-forge
• pip wheels through PyPI: pip install nvidia-nvshmem-cu12
  DeepEP is compatible with upstream NVSHMEM 3.3.9 and later.

2. Enable NVSHMEM IBGDA support Copy Section

NVSHMEM Supports two modes with different requirements. Either of the following methods can be used to enable IBGDA support.

2.1 Configure NVIDIA driver

This configuration enables traditional IBGDA support.

Modify /etc/modprobe.d/nvidia.conf:

options nvidia NVreg_EnableStreamMemOPs=1 NVreg_RegistryDwords="PeerMappingOverride=1;"

Update kernel configuration:

sudo update-initramfs -u
sudo reboot

2.2 Install GDRCopy and load the gdrdrv kernel module

This configuration enables IBGDA through asynchronous post-send operations assisted by the CPU. More information about CPU-assisted IBGDA can be found in this blog.
It comes with a small performance penalty, but can be used when modifying the driver regkeys is not an option.

Download GDRCopy
GDRCopy is available as prebuilt deb and rpm packages here. or as source code on the GDRCopy github repository.

Install GDRCopy following the instructions on the GDRCopy github repository.

Post-installation configuration Copy Chapter

When not installing NVSHMEM from RPM or deb packages, set the following environment variables in your shell configuration:

export NVSHMEM_DIR=/path/to/your/dir/to/install  # Use for DeepEP installation
export LD_LIBRARY_PATH="${NVSHMEM_DIR}/lib:$LD_LIBRARY_PATH"
export PATH="${NVSHMEM_DIR}/bin:$PATH"

Verification Copy Chapter

nvshmem-info -a # Should display details of nvshmem