# How DeepEP Enables Hook-Based Overlapping Without SM Occupation

> Discover how DeepEP enables hook-based overlapping without SM occupation by launching low-latency kernels and splitting operations into distinct phases. Learn more today!

- Repository: [DeepSeek/DeepEP](https://github.com/deepseek-ai/DeepEP)
- Tags: internals
- Published: 2026-04-25

---

**DeepEP achieves hook-based overlapping by launching low-latency kernels on the default compute stream instead of a dedicated communication stream, splitting the operation into distinct send and receive phases, and exposing a Python callback that defers the receive phase until the user explicitly invokes it.**

DeepEP is DeepSeek's high-performance communication library optimized for Mixture-of-Experts (MoE) workloads. Unlike traditional communication libraries that occupy dedicated Streaming Multiprocessor (SM) resources during asynchronous operations, DeepEP implements a unique hook-based mechanism that allows true communication-computation overlap without reserving SMs for background communication tasks.

## Compute Stream vs. Communication Stream

DeepEP normally executes low-latency all-to-all operations on a dedicated `comm_stream` to parallelize communication with computation. However, when `return_recv_hook=True` is specified, the kernel launch logic in [`csrc/deep_ep.cpp`](https://github.com/deepseek-ai/DeepEP/blob/main/csrc/deep_ep.cpp) switches the execution context to the user's current compute stream.

In [`csrc/deep_ep.cpp`](https://github.com/deepseek-ai/DeepEP/blob/main/csrc/deep_ep.cpp) (lines 1582-1585), the stream selection logic determines where the kernel executes:

```cpp
auto compute_stream = at::cuda::getCurrentCUDAStream();
auto launch_stream = return_recv_hook ? compute_stream : comm_stream;

```

By launching on `compute_stream`, the communication kernel shares SM resources with the user's existing compute kernels rather than competing for resources on a separate stream. This eliminates SM occupation because the CUDA scheduler treats the communication work as part of the normal compute workload.

## Two-Phase Kernel Architecture

The low-latency dispatch kernel is architected as a two-phase operation implemented in [`csrc/deep_ep.cpp`](https://github.com/deepseek-ai/DeepEP/blob/main/csrc/deep_ep.cpp). In hook mode, DeepEP executes only the `LOW_LATENCY_SEND_PHASE` initially, while capturing the `LOW_LATENCY_RECV_PHASE` in a callable std::function.

The phase selection and hook construction (lines 1662-1665) works as follows:

```cpp
std::optional<std::function<void()>> recv_hook = std::nullopt;
if (return_recv_hook)
    recv_hook = [=]() { launcher(LOW_LATENCY_RECV_PHASE); };

```

This design means:
- **Send phase**: Executes immediately on the compute stream, transferring data to network buffers
- **Receive phase**: Remains pending until the Python hook is invoked, allowing compute kernels to interleave

## Event Handling and Synchronization

The `low_latency_dispatch` method in [`deep_ep/buffer.py`](https://github.com/deepseek-ai/DeepEP/blob/main/deep_ep/buffer.py) returns an `EventHandle` alongside the hook. This event tracks completion of the send phase without blocking the compute stream.

The Python wrapper captures both synchronization primitives:

1. **EventHandle**: Tracks when the send phase completes on the compute stream
2. **Hook callable**: A Python-wrapped C++ lambda that launches the receive phase when called

Because the pending receive work is not actively polling on SMs (it is stored as a deferred function object), user kernels can execute on the same stream without resource contention. TheSM occupation is effectively zero until the user chooses to invoke the hook.

## Implementing Hook-Based Overlap in Practice

The following pattern from [`deep_ep/buffer.py`](https://github.com/deepseek-ai/DeepEP/blob/main/deep_ep/buffer.py) demonstrates the recommended workflow for overlapping computation with communication:

```python
import torch
from deep_ep.buffer import Buffer

# Initialize buffer and prepare tensors

# buf is a Buffer instance, x and topk_idx are prepared tensors

# 1. Dispatch with hook enabled - runs send phase only

recv_x, recv_cnt, handle, event, hook = buf.low_latency_dispatch(
    x, topk_idx,
    num_max_dispatch_tokens_per_rank=64,
    num_experts=128,
    async_finish=False,
    return_recv_hook=True  # Enable hook-based overlapping

)

# 2. Execute compute that doesn't require received data

# SM resources are fully available for user kernels here

do_other_work()

# 3. Invoke hook when received data is required

# This launches the receive phase on the compute stream

hook()
event.current_stream_wait()  # Optional: ensure completion

```

The `EventOverlap` helper class in [`deep_ep/utils.py`](https://github.com/deepseek-ai/DeepEP/blob/main/deep_ep/utils.py) provides additional utilities for managing the CUDA events that synchronize these operations across the stream boundary.

## Summary

DeepEP's hook-based overlapping mechanism eliminates SM occupation through three key design decisions:

- **Stream migration**: Communication kernels execute on the compute stream rather than a dedicated communication stream
- **Phase separation**: The send phase completes immediately while the receive phase remains deferred in a callable hook
- **Deferred execution**: No SM resources are consumed while waiting to receive; the receive phase only executes when explicitly invoked via the Python callback

This architecture allows MoE inference pipelines to maximize GPU utilization by interleaving expert computation with communication preparation without resource conflicts.

## Frequently Asked Questions

### What are the performance benefits of hook-based overlapping?

Hook-based overlapping eliminates the SM resource competition that occurs when communication kernels run on dedicated streams. By deferring the receive phase to a user-controlled hook, the GPU can dedicate 100% of SM resources to computation during the overlap window, improving throughput for compute-bound MoE layers.

### Does using return_recv_hook=True increase communication latency?

Yes, hook-based mode introduces additional latency because the receive phase cannot begin until the user invokes the hook. This mode optimizes for throughput in scenarios where computation can hide the communication latency, rather than minimizing end-to-end latency for the communication operation itself.

### When should I use hook-based overlapping versus normal async execution?

Use `return_recv_hook=True` when you have independent compute work that can execute between the send and receive phases, such as processing previous layers or preparing subsequent micro-batches. Use normal async execution (`return_recv_hook=False`) when you need minimal latency and cannot tolerate deferring the receive phase.

### How does DeepEP ensure the receive data is ready when the hook is called?

The `EventHandle` returned alongside the hook tracks CUDA stream dependencies. While the hook initiates the receive phase, the optional `event.current_stream_wait()` ensures all prior send operations complete before the received data is accessed, maintaining correct memory ordering without explicit SM occupation during the wait.