Serving MoE Models: A Deep Dive into Parallelism Strategies (TP, EP, DP)

Why this post

📚 New to MoE? If you’re not sure what a Mixture-of-Experts model is or how the router, experts, and all-to-all are actually implemented, start with the companion post Mixture-of-Experts Across the Stack — it builds the mental model from a 20-line PyTorch reference up to wide expert parallelism. This post assumes that background and goes one level up, to the serving-strategy layer. This blog post will be more practical journey on optimizing MoE serving

Serving a Mixture-of-Experts model is not the same problem as serving a dense one. A dense model has exactly two ways to spread across GPUs — tensor parallelism (TP) and data parallelism (DP). An MoE adds a third axis, expert parallelism (EP). This post is a deep dive into those three strategies as they apply to MoE serving — what each shards, what collective it puts on the critical path, what it costs in memory, and when to pick it. It’s grounded in three places: the vLLM source (traced down to the exact lines that pick the kernel), a Qwen3-235B-A22B reward-judge benchmark I ran on 8×H200, and three references that corroborate each other — AMD’s vLLM MoE parallelism guide, SGLang’s large-scale EP report, and the vLLM data-parallel docs. The MoE mechanics themselves (router, grouped GEMM, the all-to-all kernel) get a full treatment in the companion MoE post; here we’re one level up, at the serving-strategy layer.

---

Benchmark setup

Every measured number in this post comes from one rig — worth stating once so the figures (55 GiB/GPU, ~72 req/s, the traces) have context.

Model — Qwen3-235B-A22B-Instruct-2507. A sparse MoE: 235 B total parameters, 22 B active per token (the “A22B”).

layers	94
hidden size	4096
attention	64 query heads / 4 KV heads (GQA), head_dim 128
experts	128 routed, top-8 per token (+ shared expert)
MoE intermediate	1536 per expert
vocab	151,936
dtype	bf16 (≈ 470 GB of weights)

Hardware. One node, 8× NVIDIA H200 (141 GB HBM3e each), NVLink/NVSwitch intra-node. All parallelism in this post is within this single node — which is exactly why TP’s intra-node all-reduce stays cheap and why the PP experiment below is a single-node curiosity, not a deployment.

Workload — the reward judge. A driver replays real reward-model requests against the vLLM server, round-robin across replicas (mirroring verl’s reward router), closed-loop at fixed concurrency:

prompt length	~2,000–4,700 tokens
output length	4 tokens (the judge emits a short verdict; some calls also request top-20 logprobs)
concurrency	64 in-flight
dataset	2,048 real requests
in the RL loop	2,048 reward calls per training step
API	/v1/completions, client-templated prompt, temperature=0

That shape — long prompt, 4-token output — is the crux: it makes the reward prefill-compute-bound, so the parallelism choice is decided by prefill GEMM efficiency + memory, not decode bandwidth. It’s also why prefix caching dominates — the ~2,500-token job posting is shared across a member’s 4 candidate summaries, so warm throughput (~72 req/s at TP8+EP) is ~12× the cold first pass (~6 req/s).

The baseline: Tensor Parallelism (TP=8)

TP shards everything and pays every layer. Each rank holds 1/TP of every weight matrix — attention q/k/v/o, the experts, the lm_head. After each sharded matmul it must all-reduce the partial activations. Lowest per-GPU memory, but a collective on the critical path of every layer.

For MoE specifically, TP without EP means every expert is tensor-sharded across all ranks. The router still picks the top-k experts for each token, but those chosen experts are not owned by one GPU. Each GPU holds a slice of every expert’s weights, runs its slice of the expert GEMM, and then participates in the collective that merges the partial outputs. So TP saves memory by slicing the whole model evenly, but it also keeps the MoE layer coupled across every TP rank.

Tensor parallelism without expert parallelism: each expert is split across the TP group, so one routed expert still involves all ranks. Figure from AMD’s vLLM MoE parallelism guide.

profiling result

_TP8 without EP profiling overview: the execute_context_87(111606)_generation_112(112) means 87 context reqeust with 111606 prefill token, and 112 decode request.

Detailed TP8 without EP profiling breakdown: compute remains the largest bucket, with all-reduce as the main communication cost and no all-to-all path.

We can check the overhead of communication by analyzing the trace profiling:

	% of GPU-busy time
compute (attention + MoE GEMMs)	~75%
all-reduce	~20%
memcpy	~5%

It’s essentially all ncclDevKernel_AllReduce_Sum — the two all-reduces per layer (after the attention O-proj, after the MoE down-proj) per-rank work is balanced, so the all-reduce is near-pure transfer with negligible barrier wait — a flat ~20% on every rank (unlike EP, where expert-load skew inflates it; see §2).

request througput: 40 req/s

🔧 Why it’s one all-reduce per block, not one per matmul?

“All-reduce after each sharded matmul” is the right intuition but counts slightly too many: an FFN/expert has two matmuls and an activation, yet pays only one all-reduce. The trick is the Megatron column-parallel → element-wise → row-parallel pairing:

• Up-proj (W_gate/W_up) is column-parallel — sharded along the intermediate dim I. Rank i computes a disjoint slice of the intermediate, [N, I/P]. That’s a concatenation, not a partial sum — each rank’s neurons are already complete, so there is nothing to reduce after the first GEMM.
• SwiGLU is element-wise along I: silu(gate_i) ⊙ up_i runs entirely within rank i’s slice (gate and up are sharded identically), so the activation never needs another rank’s data. The intermediate stays sharded straight through.
• Down-proj (W_down) is row-parallel — now every rank produces a partial sum of the full output [N, H]. That’s what the single all-reduce sums.

So you pay a collective only when row-parallel turns the sharded intermediate back into a sum. If the middle op mixed across I (a norm or softmax over the intermediate dim), you’d be forced to all-gather/all-reduce before it — but element-wise activations never do, by design. Net: attention = 1 all-reduce (after the O-proj), FFN/MoE = 1 all-reduce (after the down-proj). (In the EP case the expert’s two matmuls + SwiGLU are all local to one rank — zero intra-expert comm — and the one all-reduce is the cross-rank expert combine instead.)

---

Tensor Parallel + Expert Parallelism

--enable-expert-parallel changes how the expert weights are split — from “shard every expert across ranks” (TP-style) to “give each rank a few whole experts” (EP-style) — and changes the MoE communication accordingly. Attention is still parallelized by whatever TP/DP you set.

Tensor parallelism with expert parallelism: attention and dense weights still follow TP, while MoE expert weights are partitioned as whole experts across the ranks. Figure from AMD’s vLLM MoE parallelism guide.

Profiling Result

_TP8 + EP8 (dp=1) profiling timeline: MoE experts are owned as whole experts, but the visible communication is still NCCL all-reduce, and we can see there are some imbalance in the serving

request througput: 44 req/s (not too much change since we are not saving any compute/communication overlap)

Data Parallel + Expert Parallelism

Every config above ran dp=1, so the MoE combine was an all-reduce and the all-to-all never appeared. To actually see it I booted the judge as DP8 × EP8 × TP1 — eight data-parallel attention replicas, experts expert-parallel across all eight (vllm serve … --data-parallel-size 8 --enable-expert-parallel). This is the AMD guide’s DP-with-EP layout.

Data parallelism with expert parallelism: the dense path is replicated across DP ranks, while routed tokens cross ranks to reach the experts they select. Figure from AMD’s vLLM MoE parallelism guide.

Profiling Result

DP8 + EP8 profiling timeline: the token count per rank is small because we have to set max_num_batched_tokens 131072→16384 to make sure it does not OOM. The all-gather/reduce-scatter path spends a large fraction of time in NCCL broadcast and reduce kernels instead of useful expert GEMM.

kernel (rank 7)	self-CUDA %	role
fused_moe	43%	expert grouped-GEMM (overlaps comm)
ncclDevKernel_Broadcast_RING_LL	31.5%	dispatch (tokens → expert ranks)
ncclDevKernel_Reduce_Sum_bf16_RING_LL	15.9%	combine (expert outputs → token ranks)

request througput: 20 req/s

Communication is ~47% of GPU-busy time — vs ~20% all-reduce for TP8+EP8 — which is why one-node DP+EP was ~2× slower in my e2e judge (reward phase 99 s vs 46 s). But why broadcast + reduce-scatter instead of one all-to-all kernel?

vLLM’s default single-node all2all backend is AgRsAll2AllManager — “AllGather / ReduceScatter”. The serve log says it outright: Using AgRsAll2AllManager all2all. Its dispatch is an all_gatherv and its combine is a reduce_scatterv:

class AgRsAll2AllManager(All2AllManagerBase):
    def dispatch(self, hidden_states, topk_weights, topk_ids, ...):
        sizes = dp_metadata.get_chunk_sizes_across_dp_rank()      # tokens per DP rank
        gathered = dist_group.all_gatherv([hidden_states, ...], dim=0, sizes=sizes)
        ...      # every rank now holds ALL tokens, runs its local experts on them
    def combine(self, hidden_states, ...):
        sizes = dp_metadata.get_chunk_sizes_across_dp_rank()
        return dist_group.reduce_scatterv(hidden_states, dim=0, sizes=sizes)

So the “all-to-all” is realized as: all-gather every rank’s tokens (dispatch) → each rank runs its local experts over the full token set → reduce-scatter the partial outputs back (combine). That’s the reduce-scatter half. The broadcasts come from the all-gather. The router is load-imbalanced, so each DP rank holds a different number of tokens — this is a variable-size all-gather (all_gatherv), and NCCL has no symmetric primitive for ragged sizes. vLLM’s pynccl.all_gatherv builds it from P broadcasts — one per DP rank, each rank as root for its own ragged chunk — but wraps them in an NCCL group:

def all_gatherv(self, output_tensor, input_tensor, sizes, ...):
    self.nccl.ncclGroupStart()
    for root, split_size in enumerate(sizes):       # P broadcast OPS, one per DP rank
        dst_slice = output_tensor[split_offset : split_offset + split_size]
        self.nccl.ncclBroadcast(input_tensor, dst_slice, ..., root, self.comm, ...)
        split_offset += split_size
    self.nccl.ncclGroupEnd()                        # ...fused into ONE kernel launch

Don’t be fooled into counting eight kernels — ncclGroupStart()/ncclGroupEnd() is exactly NCCL’s mechanism for aggregating the enclosed ops into a single fused device-kernel launch. The trace confirms it: per rank, ncclDevKernel_Broadcast_RING_LL fires 4,418 times — identical to vllm::moe_forward’s 4,418 (and to the reduce-scatter’s 4,418). So it’s one broadcast kernel per MoE-layer dispatch, not eight; the P=8 is the number of roots stitched into that one kernel (the ragged all-gather), not the number of launches. The “lots of broadcast” you see in the timeline is one-per-layer × 94 layers × every decode/prefill step — plenty of launches, but each is a single fused ragged all-gather, not an 8× loop of separate collectives.

The “sometimes all-gather + reduce-scatter” is the same code taking its balanced fast-path. cuda_communicator.all_gatherv checks whether the per-rank sizes are equal:

if sizes is not None and all(s == sizes[0] for s in sizes):
    sizes = None                       # balanced -> drop the ragged sizes
...
if sizes is not None:
    pynccl_comm.all_gatherv(...)       # imbalanced -> the broadcast loop above
else:
    pynccl_comm.all_gather(...)        # balanced  -> one symmetric ncclAllGather

When the router happens to give every DP rank the same token count that step, sizes collapses to None and you get a single ncclAllGather; when it doesn’t (the usual case), you fall back to the broadcast loop. The combine has the identical split — reduce_scatterv for imbalanced, reduce_scatter for balanced. So the trace’s flip-flop between broadcast and all-gather + reduce-scatter is a per-step readout of how balanced the router’s token distribution was that iteration — broadcasts when skewed, a clean all-gather when even.

⚠️ This is the naive backend, not DeepEP. AgRs all-gathers the full token set to every rank; DeepEP replaces dispatch with a true point-to-point all-to-all that sends each token only to the ranks owning its top-k experts (≈ 5.25 of 8, not 8). Without DeepEP built on the box, one-node DP+EP moves far more bytes than TP8+EP8’s NVLS all-reduce and loses. The AMD-blog DP+EP win assumes the efficient kernel and enough scale (EP ≫ k over a slow inter-node fabric) to amortize it.

---

Communication Overhead Comparison

Here is the fact that cost me an afternoon of confusion. In vLLM, the all-to-all only exists when there is data parallelism. From fused_moe/config.py:

@property
def use_all2all_kernels(self):
    return self.dp_size > 1 and self.use_ep

Every dispatch/combine backend — DeepEP, allgather_reducescatter, the naive path — is gated on this flag. If it’s False, vLLM’s maybe_make_prepare_finalize factory returns None, no dispatch/combine kernel is built, and the MoE falls back to the standard grouped-GEMM whose partial outputs are merged with a plain all-reduce. So:

Config	MoE communication
TP only (dp=1), EP off	all-reduce
TP + EP (dp=1)	all-reduce
DP, EP off	independent replicas (allgather/reducescatter inside a replica’s MoE)
DP + EP (dp>1), DeepEP	all-to-all — point-to-point dispatch + combine
DP + EP (dp>1), naive allgather_reducescatter (vLLM’s single-node default)	all-gather (dispatch) + reduce-scatter (combine)

That last row is the one that bit me. The flag flips on the “all-to-all path,” but on one node vLLM doesn’t run a true token-routing all-to-all — it runs AgRsAll2AllManager: all-gather every rank’s tokens to every rank, then reduce-scatter the expert outputs back. The “broadcasts” you see in the trace (§3) are that all-gather — a ragged all_gatherv that NCCL expresses as P broadcasts fused into one kernel. A real all-to-all is reserved for DeepEP, and that distinction changes the byte count below.

The communication bytes, calculated

The table above says what collective each config runs; this says how much it moves. Take one layer’s activation tensor as the unit: S = N·H·2 bytes (bf16) — for the judge’s prefill chunk (N=32768, H=4096) that’s S = 256 MiB. The primitives:

• a ring all-reduce moves 2(P−1)·S across the group (P=8 → 14·S);
• the ideal MoE all-to-all (DeepEP) moves 2·fanout·S, where fanout = P·[1−(1−1/P)^k] ≈ 5.25 for top-k=8 on 8 ranks (each token reaches ~5.25 of 8 ranks, not all) → 10.5·S;
• but the broadcast-backed all-gather that vLLM actually runs (AgRs, the default) moves the full set: an all-gather of S costs (P−1)·S, the reduce-scatter combine another (P−1)·S, so dispatch+combine = 2(P−1)·S = 14·S — byte-for-byte an all-reduce, and ~33% above the idealized all-to-all (14·S vs 10.5·S). The “fanout” is effectively the full P=8 (every token’s hidden state is broadcast to every rank), not 5.25, because nobody pruned the routing before moving the data.
• a PP hand-off is one point-to-point send of S at each of the (p−1) stage boundaries.

Over L=94 layers:

config (8 GPUs)	per-layer comm	bytes / forward / GPU	aggregate / forward	dominant cost
TP8 · EP8 (dp=1)	2× all-reduce (attn + MoE)	82 GiB	658 GiB	per-layer all-reduce
TP8 · EP off (dp=1, TP-MoE)	2× all-reduce (attn + MoE)	82 GiB (identical)	658 GiB	same comm; EP only moves compute/memory
TP1 · EP8 · DP8, AgRs (actual)	all-gather + reduce-scatter (= 14·S, no attn all-reduce)	41 GiB	329 GiB	broadcast all-gather + §4 lockstep
↳ if DeepEP (true all-to-all)	2× all-to-all = 10.5·S	31 GiB	247 GiB	a real routing-aware dispatch

(Per-layer in S units: all-reduce configs = 28·S/layer aggregate = 2 × 1.75·S per GPU; AgRs DP+EP = 2(P−1)·S = 14·S/layer = 1.75·S per GPU — same per-GPU as one all-reduce; the DeepEP ideal would be 2·fanout·S = 10.5·S; PP = (p−1)·S for the whole forward.)

Three reads:

• TP8·EP8 and TP8·EP-off move the exact same bytes. Both all-reduce twice per layer over [N,H]. EP changes which matmul each rank runs (whole experts vs sliced) and the weight memory (55 vs 110 GiB/GPU) — not the communication. The ~4 s/step gap between them (§6) is a GEMM-shape effect, not a comm one.
• TP1·EP8·DP8 (AgRs) moves half TP8’s bytes (41 vs 82 GiB/GPU) yet is ~2× slower — and the trace agrees (47% comm vs 20%). Two reasons the byte count misleads. (1) It’s the wrong kind of collective. TP’s all-reduce runs over NVLink with NVLS / in-switch reduction (SHARP) — the NVSwitch sums the partials in the network, so the reduction is “free” and effective bandwidth far exceeds the wire. AgRs’s dispatch is a ragged all-gather built from broadcasts: pure data motion, no in-switch reduction, and a broadcast tree is less bandwidth-optimal than the ring/NVLS all-reduce — so each of those 41 GiB crawls compared to TP’s 82 GiB. (2) It’s gated by the slowest rank plus a per-step barrier. The all-gather waits on whichever rank holds the most tokens (I measured 1.95× max/min imbalance), and every forward pays the CPU all-reduce that agrees on DP padding (§4) — neither tax exists for the balanced, barrier-free TP all-reduce. Net: fewer bytes, moved far slower, plus lockstep (see appendix). Swapping in DeepEP cuts the bytes (41→31 GiB) and replaces the broadcast with a real point-to-point dispatch — but on one node at k=8 even that loses to NVLS all-reduce; DP+EP only pays once EP ≫ k (sparser routing, more ranks) over a slow inter-node fabric.
• PP moves almost nothing — a few hundred MiB of boundary p2p against tens of GiB of per-layer collectives. Its cost isn’t bytes, it’s the bubble; it earns its place offline / cross-node, where a deep request backlog hides the bubble and a small p2p beats a cross-node all-reduce.

Memory consumption

Before KV cache, activations, CUDA graphs, and temporary all-to-all buffers, the theoretical floor is just bf16 weights. For Qwen3-235B-A22B, the expert weights dominate:

• per expert per layer: 3 × hidden × intermediate = 3 × 4096 × 1536 = 18.9M params
• routed experts: 18.9M × 128 × 94 ≈ 227B expert params
• non-expert remainder: 235B - 227B ≈ 8B params (attention, norms, router, embeddings/lm head, shared pieces)

So the useful mental model is:

\[ext{weights/GPU} \approx 2 \cdot \left(\frac{P_{dense}}{TP} + \frac{P_{expert}}{EP}\right)\]

where bf16 is 2 bytes/param, TP shards the dense path inside each DP replica, and EP shards the routed experts as whole experts. For this 8-GPU node:

layout	dense params/GPU	expert params/GPU	theoretical weights/GPU	what it means
TP8, EP off	8B / 8	227B / 8	~54.7 GiB	every layer, including every expert, is tensor-sharded
TP8 + EP8 (dp=1)	8B / 8	227B / 8	~54.7 GiB	same byte floor as TP8, but experts are whole-expert sharded
DP2 / TP4 + EP8	8B / 4	227B / 8	~56.5 GiB	dense path is less sharded because there are two DP replicas
DP4 / TP2 + EP8	8B / 2	227B / 8	~60.3 GiB	more DP attention, more replicated dense weight
DP8 / TP1 + EP8	8B	227B / 8	~67.6 GiB	experts are still sharded, but dense weights are fully replicated
DP8, EP off	235B	included	~437.7 GiB	impossible on one H200; full model replica per GPU

Two takeaways matter for this benchmark. First, TP8 and TP8+EP8 do not materially differ in theoretical weight bytes: both spread the 235B parameters across the same 8 GPUs. EP changes the shape of the local work — whole experts instead of slices of every expert — not the total weight floor. Second, DP attention spends memory by replicating the dense remainder. That ~8B dense tail is small compared with the experts, but every step from TP8 toward DP8 adds a few GiB/GPU before KV cache and runtime buffers, which is why the measured peak can move from comfortable to near-OOM even when the bf16 weight table looks close.

Appendix: How the layer, runner, and kernel get wired up at init

The all-reduce-vs-all-to-all decision above is really a kernel-selection decision made once, at model build time. It’s worth seeing the whole init pipeline, because vLLM 0.22 splits the MoE into three objects — a layer (FusedMoE), an execution runner (MoERunner), and a kernel (FusedMoEKernel = a prepare/finalize + an experts GEMM) — and “which kernel runs” is decided by which dp_size/quant combination you booted with.

Step 1 — the layer picks a quant method and allocates the (sharded) weights. FusedMoE.__init__ first resolves a quant method in _get_quant_method: self.quant_config.get_quant_method(...) returns e.g. Fp8MoEMethod / CompressedTensorsWNA16MoEMethod, and if there’s no quant config it falls back to UnquantizedFusedMoEMethod. That method’s create_weights allocates the expert weight tensors — and because EP shards by whole expert, each rank only allocates its local_num_experts (16 of 128 for our TP8+EP8 judge), which is why weights land at ~55 GiB/GPU instead of 110.

Step 2 — the layer constructs the runner. self.runner = MoERunner(...) hands the quant method, router, gate, and shared experts to MoERunner.__init__. The runner is the execution object — it wraps the shared experts in a SharedExperts helper and calls _select_forward, which binds the forward entry to the registered custom op torch.ops.vllm.moe_forward (or the _shared variant when shared experts exist). Note the runner does not itself choose a GEMM — it just holds self._quant_method and later calls through it.

Step 3 — the kernel is chosen after weights load, by pairing a dispatch impl with an experts GEMM. This is the crux. maybe_init_modular_kernel assembles a modular kernel from two halves:

• the prepare/finalize (dispatch/combine) object from maybe_make_prepare_finalize — None at dp=1, as traced above;
• the experts GEMM from select_gemm_impl(prepare_finalize, layer), which each quant method overrides to return the right FusedMoEExpertsModular (Triton / DeepGEMM / CUTLASS / TRT-LLM / …) based on the all2all impl and dtype.

FusedMoEModularMethod.make glues them together — note the experts GEMM is literally a function of the dispatch object:

return FusedMoEModularMethod(
    old_quant_method,
    FusedMoEKernel(
        prepare_finalize,                                       # dispatch / combine
        old_quant_method.select_gemm_impl(prepare_finalize, moe_layer),  # experts GEMM, chosen from it
        inplace=inplace,
    ),
)

The layer then swaps self.quant_method to this modular method. But at dp=1 the first half is None, so maybe_init_modular_kernel short-circuits (the if prepare_finalize is not None guard) and never builds a modular kernel — the quant method keeps its monolithic path (is_monolithic = True), i.e. a single fused-experts GEMM with no cross-rank dispatch, whose partial output is reduced by the runner’s all-reduce. So “EP8/TP8/dp=1 → all-reduce” is exactly “no prepare/finalize was built, so the monolithic kernel was kept.”

Step 4 — at runtime it’s a thin call-through. FusedMoE.forward calls runner.forward, which routes through _apply_quant_method: is_monolithic ? apply_monolithic : apply. The monolithic branch is the one our judge takes; its output then hits _maybe_reduce_final_output and the all-reduce fires. The modular branch (the apply side) is where DeepEP/CUTLASS dispatch-combine would have run, had dp_size > 1 selected it at Step 3.

So the whole “which collective” question is resolved at build time, in Step 3, by whether maybe_make_prepare_finalize returned a dispatch object — everything downstream is just executing the kernel that selection already locked in.

Appendix: The hidden cost of DP+EP: lockstep

Plain DP ranks are embarrassingly parallel. Add EP and they become a single lockstepped machine, because the expert all-to-all is a collective — every rank must reach it together, every MoE layer.

Three concrete taxes, all of which I saw or the docs name:

1. A per-step barrier. Routing is dynamic, so before each all-to-all the ranks must agree on shapes. vLLM logs it: Using CPU all reduce to synchronize DP padding between ranks — a tiny CPU all-reduce every forward to take the max token count and pad to it.
2. Dummy forwards for idle ranks. The vLLM docs: “expert layers across all ranks are required to synchronize during every forward pass, even when there are fewer requests than DP ranks” — so an idle rank still runs an empty forward (managed by a “separate DP Coordinator process”) just to keep the collective alive.
3. Straggler tail. Everyone waits for the slowest/longest-prompt rank each step.

What the dummy step (tax #2) looks like in the profile — and where it’s triggered

You can see it directly. Filter the trace to kernels with no execute_context parent: that annotation is built only inside the real forward (gpu_worker.py:759 — the "execute_context_" + num_ctx_requests … string wrapped around execute_model), so anything outside it is not a real step — it’s a _dummy_run. On an idle DP rank, that region is a full MoE layer with zero attention:

kernel (one dummy layer)	present?
ncclDevKernel_Broadcast_RING_LL (all-gather dispatch)	✅
topkGating → fused_moe_kernel → ncclDevKernel_Reduce_Sum (combine)	✅
up/down-proj GEMMs, silu SwiGLU, RMSNorm, reshape_and_cache_flash	✅ (cheap per-layer plumbing)
flash::FlashAttnFwd… (attention compute)	❌ absent

The rank has no real sequences, so attention has nothing to compute — it’s just turning the MoE all-gather/reduce-scatter crank to stay lockstepped (attention is also a cudagraph splitting op, splitting_ops: ['vllm::unified_attention_with_output', …], so it sits outside the captured MoE pieces regardless). The trigger is the per-step DP sync coordinate_batch_across_dp: every forward, the ranks CPU-all-reduce their token counts (tax #1); if any rank has work, an idle rank calls self._dummy_run(1) — “dummy run to ensure coordinate_batch_across_dp is called … to avoid out of sync issues” — instead of skipping the collective. That’s GPU time spent moving and computing the MoE purely to keep an idle rank in step: the lockstep tax, with a kernel signature you can grep for.

This is why “wide EP” is a systems project, not a flag. The state of the art exists to claw the lockstep back:

• Two-Batch Overlap (TBO) — split the batch so one micro-batch’s expert compute hides the other’s all-to-all. SGLang reports 27–35% prefill throughput from it — but −27% at 32 tokens/device, i.e. it hurts at small batch.
• EPLB (expert-parallel load balancing) — replicate hot experts as redundant copies. SGLang: 1.49× prefill / 2.54× decode.
• A real all-to-all kernel — DeepEP, which moves only the routed top-k tokens over NVSHMEM (see the NVSHMEM deep dive for the one-sided transport it rides on). Without DeepEP, vLLM’s default allgather_reducescatter gathers every rank’s tokens to every rank — which in my benchmark made the MoE GEMM time go up (595 → 905 ms) and total comm go up (562 → ~893 ms) versus the TP8 all-reduce. The fallback is a regression; the win needs the real kernel.

---

5. When to choose what

Putting it together — a decision framework, with the crossover numbers from AMD’s benchmarks:

By concurrency (the first cut)

• Latency / low concurrency (≤128 in-flight): TP. One replica, every GPU on every token, lowest TTFT. AMD measured “40–86% higher throughput” and “80% lower TTFT” for TP vs DP here; for DeepSeek-R1, TP8+EP was “52% higher throughput” than DP8+EP at low load.
• Throughput / high concurrency (≥512 in-flight): DP (replicas). AMD: “16–47% higher throughput at scale”; DeepSeek-R1 DP8+EP hit “7,114 tok/s (47% higher than TP8+EP).”
• Crossover: 256–512 concurrent requests across the models they tested. Below it TP wins, above it DP wins.

By the MoE’s footprint and load (the second cut)

(A dense model has no experts and thus no EP — it’s just TP for latency / a model too big for one GPU, DP replicas for throughput. Everything below is MoE-specific, which is where the real choices live.)

• MoE that fits in ≤8 GPUs, single replica, latency- or memory-bound: TP + EP (dp=1). Experts sharded for the memory win, combined with an all-reduce — vLLM’s weight-only fold (§2 aside). Simple, no lockstep, no DeepEP. (This is the reward judge in §6.)
• MoE at scale / decode-heavy / throughput-bound: DP attention + EP all-to-all, with DeepEP + EPLB + TBO. This is the DeepSeek/SGLang wide-EP regime — and the only vLLM path that actually builds the dispatch/combine kernel. SGLang ran EP32 for prefill, EP72 for decode (disaggregated), hitting $0.20 / 1M output tokens and “up to 5×” over a TP16 baseline.

Two routing rules that override the above

• Expert activation density. Ultra-sparse models can lose to EP because the all-to-all dominates: AMD found Llama-4-Maverick (0.78% density) runs 7–12% faster with EP off. Their rule of thumb: < 1% density → EP off; > 3% → EP on.
• MLA / single-KV-head models (DeepSeek): DP attention is essentially required — a single latent KV head can’t be sharded usefully by TP, so you replicate attention (DP) and EP the experts.

The one-paragraph summary

If you take one thing away: on a single node, the only knob that changes how much data moves is TP’s all-reduce — and “expert parallel” doesn’t touch it. Tensor parallelism shards every weight matrix and pays one all-reduce per block (~20% of the step, flat across ranks); adding --enable-expert-parallel at dp=1 merely swaps sliced experts for whole experts — a memory-and-GEMM-shape change, not a communication one, since the partial outputs are still merged by that same all-reduce. The token-routing all-to-all everyone associates with EP exists only at dp>1 (use_all2all_kernels = dp_size > 1 and use_ep), and even then vLLM’s single-node default isn’t a real all-to-all: it’s allgather_reducescatter — broadcast every rank’s tokens to every rank, reduce-scatter the outputs back — which moves a full all-reduce’s worth of bytes (2(P−1)·S = 14·S, ~41 GiB/GPU), ~33% above the idealized DeepEP all-to-all (10.5·S) and ~2× slower than TP despite moving fewer total bytes, because the broadcast gets no in-switch NVLS reduction, is gated by the hottest (1.95×-imbalanced) rank, and drags a per-forward lockstep barrier (§4). What DP genuinely buys is 8× KV-cache capacity — each rank caches its own requests — and the throughput headroom that follows, which is why the crossover is concurrency: TP below ~256–512 in-flight, DP above. For my prefill-bound reward judge on 8×H200 that verdict was unambiguous — TP8+EP8 wins outright; DP+EP only earns its lockstep once you’re multi-node, decode-heavy, and running DeepEP + EPLB to claw the all-to-all back. The flag is one line; making it pay is a systems project.

---

References: AMD ROCm — vLLM MoE parallelism guide, SGLang — Large-Scale Expert Parallelism, vLLM — Data Parallel Deployment docs, NVIDIA Megatron-LM. Source traced against vLLM v0.22.0 and Megatron-core core_v0.17.1 (line anchors pinned to those tags); benchmark numbers from a Qwen3-235B-A22B reward judge on 8×H200 (vLLM 0.18/0.22). Companion posts: MoE across the stack, gpu_memory_utilization, NVSHMEM. Frameworks move fast — re-pin before relying on a line number.