Explaining DeepEP Dispatch (Tourist's Guide to Expert-Parallel Libraries, Pt. 1)

June 30, 2026

There are plenty of blog posts about MoEs. What’s rarer is a step-by-step walkthrough of the communication pattern underneath expert parallelism, and the kernel ideas that make it fast. This post tries to fill that gap — MoEs from first principles, then DeepEP’s dispatch kernel.

Assumes familiarity with Transformers and some GPU background.

Background: the MoE computation

An MoE layer replaces the feed-forward block in a transformer. The pitch is familiar: similar quality with far fewer parameters activated per token.

Compared to a plain FFN (matmul → activation → matmul), an MoE layer adds a few moving parts:

• Routing. Each token goes through a learned router that picks which expert(s) should process it.
• Sharded experts. Experts live across ranks/GPUs. In the diagram, each rank holds two experts — but you can put any number on a GPU, or even shard one expert across multiple GPUs.
• Dispatch. Tokens have to move from the rank that owns them to the rank that owns their assigned expert. That’s the central communication problem.
• Combine. A token can be routed to multiple experts. After those experts run, you combine their outputs — typically a weighted sum using the router scores.

Naive expert parallelism

People often say MoEs are an “all-to-all” workload. Here’s what that actually means.

An all-to-all is like a transpose over the message and processor dimensions: every rank sends a chunk to every other rank. What the classic diagram usually doesn’t show is that those chunks don’t have to be the same size. The triangle can be larger than the circle; you just declare each chunk’s length.

That variable-size pattern is exactly what MoEs need. Each expert is tied to a GPU, and each expert gets some number of tokens — so each GPU needs a different number of tokens from every other GPU.

In the naive version, you:

1. Count tokens per expert / per GPU.
2. Pack them into a tensor of the right shape.
3. Launch an all-to-all.

Those all-to-alls are dedicated kernels: the CPU launches GPU kernels that do the peer-to-peer transfers that implement the collective. Before DeepEP, that was the default paradigm — CPU-driven all-to-all everywhere. It works, but you pay packing overhead and kernel-launch overhead.

DeepEP and the explosion of EP libraries

DeepEP’s motivation is straightforward: instead of launching an all-to-all from the CPU, do device-initiated communication with NVSHMEM + IBGDA.

Under the hood, DeepEP is a dispatch / combine kernel on top of NVSHMEM — with careful design choices like warp specialization and hot-expert handling.

Design of the DeepEP dispatch kernel

Pseudocode:

__global__ moe_dispatch_kernel() {
    int my_thread_idx = threadIdx.x;
    int warp_id = my_thread_idx / 32;
    int lane_id = my_thread_idx % 32;

    // Shared across warps in the block:
    // [num_toks_to_rank_0, num_toks_to_rank_1, ...]
    __shared__ int shared_token_counts[MAX_RANKS];
    if (my_thread_idx < MAX_RANKS) {
        shared_token_counts[my_thread_idx] = 0;
    }

    __syncthreads();

    // --------------------------------------------------------------------
    // PHASE 1: WORKLOAD (both warps active concurrently)
    // --------------------------------------------------------------------
    if (warp_id == COORDINATOR_WARP_ID) {
        // Count tokens per destination rank via shared-memory atomics.
        // Build the map: {r0 gets 2 toks, r1 gets 1 tok, ...}
        for (int i = lane_id; i < num_tokens; i += 32) {
            int target_rank = token_to_rank[i];
            atomicAdd(&shared_token_counts[target_rank], 1);
        }

    } else if (warp_id == SENDER_WARP_ID) {
        // Actually put the token data on the wire.
        // Stride so each lane sends a share of the tokens.
        for (int i = lane_id; i < num_tokens; i += 32) {
            int target_rank = token_to_rank[i];
            nvshmem_put_token_nbi(globalTokens[i], target_rank);  // async put
        }
        nvshmem_quiet();  // wait until outbound network is clear
    }

    // Block-wide barrier: coordinator finished counting;
    // sender finished putting tokens and quieted.
    __syncthreads();

    // --------------------------------------------------------------------
    // PHASE 2: SIGNAL PEERS
    // Peer ranks still need to know how many tokens landed in their memory.
    // --------------------------------------------------------------------
    // Only the coordinator warp signals; lane 0 does the work.
    if (warp_id == COORDINATOR_WARP_ID && lane_id == 0) {
        for (int peer_rank = 0; peer_rank < MAX_RANKS; peer_rank++) {
            int total_sent = shared_token_counts[peer_rank];

            if (total_sent > 0) {
                // Ring the doorbell on the destination GPU
                nvshmem_atomic_add(remote_sync_counters, total_sent, peer_rank);
            }
        }

        // Flush so the signal clears PCIe / NIC
        nvshmem_quiet();
    }
}

Important note on nvshmem_quiet: when nvshmem_quiet returns, the fabric has delivered your data into the receiver’s memory — bytes crossed the network, hit the destination NIC, and are committed to remote VRAM, visible to compute cores on that GPU. The receiver still doesn’t know the data arrived.

NVSHMEM is one-sided: the receiver needs an explicit signal, flag, or barrier before it’s allowed to read the completed payload.

Translation

Say I’m rank 0. I launch one block with two warps (64 threads).

• Warp 0 is the sender; warp 1 is the coordinator.
• Sender: each lane sends its share of tokens to their target ranks (non-blocking), then calls nvshmem_quiet() and waits until the data has landed.
• Both warps hit __syncthreads(). Sender has put everything in flight and quieted; coordinator has finished the token counts.
• Coordinator: one thread nvshmem_atomic_adds the token counts onto peer ranks, then quiet()s so the signal is guaranteed delivered.
• Receiver spin-waits until its local counter is nonzero:

while (nvshmem_atomic_fetch(local_counter) == 0) {
    // spin
}

Then it processes the arrived data.

TO BE CONTINUED: DeepEP Combine, PPLX-Kernels, etc.