Mixture of Experts

Here is the symptom, in bytes. Take one DeepSeek-V3 MoE layer on one H100 and count the routed expert weights it must pull out of HBM for one decode step. At batch 1 the router picks 8 experts of 256, so the kernel reads 8 experts' worth: $8 \times 3 \times 7168 \times 2048 = 352.3$ M fp8 bytes. At batch 256 — a modest serving batch — 2,048 independent top-8 draws land on essentially every expert, and the kernel reads all 256: 11.27 GB. 256× the tokens, 32× the bytes.

Now do the same for a dense layer of the same width. Batch 1 reads the weight matrix once. Batch 256 reads the weight matrix once. 256× the tokens, 1× the bytes.

That is the whole chapter in three numbers. A dense decode GEMM amortises its weight read across the batch, so its arithmetic intensity climbs at $2B/b$ and hits the roofline ridge quickly (§0.4). An MoE layer's climbs at $2Bk/(bE)$ once the batch hits every expert — a factor $E/k$ slower, 32 for DeepSeek-V3. MoE looks like a compute optimisation. It is a memory-capacity problem wearing a compute problem's clothes.

§5.3 owns the answer at cluster scale — expert parallelism, the all-to-all, EPLB. This chapter owns MoE as a model and a kernel on one device: what the router computes, why the expert GEMM cannot be a GEMM, and where the bytes go.

A dense feed-forward block is one matrix applied to every token. An MoE block is $E$ matrices, a small learned classifier that picks $k$ of them per token, and a weighted sum of the $k$ results. Every token still does $k$ matrices' worth of arithmetic — the FLOPs per token are fixed and small. But which $k$ is decided at run time, per token, and the union over a batch grows fast. The consequence: the layer's compute scales with $Bk$ and its traffic scales with the number of distinct experts the batch touched, which saturates at $E$.

Mechanically, that turns one clean GEMM into a scatter, a set of ragged GEMMs, and a gather.

Two things in that picture are worth pausing on. First, the alignment step: the grouped GEMM needs each expert's row count to be a multiple of the tile height, so 512 real rows become about 3,552 padded ones. Second, the weight traffic: 64 tokens touched 87% of the layer's experts. There is no small-batch regime in which MoE reads only a little.

Shapes, sourced

vLLM checks in a table of real MoE shapes for its kernel benchmark. Two of its rows are the two models this chapter uses:

# Model configurations: (name, E, N, K, topk, dtype_str, use_fp8, block_shape)
# N = moe_intermediate_size // tp_size (the value used in config file lookup)
MODELS = [
# ...
    # Qwen3-30B-A3B: E=128, topk=8, N=768, K=2048
    ("Qwen3-MoE bf16", 128, 768, 2048, 8, None, False, None),
    # DeepSeek-V3 / MiMo-V2-Flash: E=256, topk=8, N=2048, K=7168
    ("DeepSeek-V3 bf16", 256, 2048, 7168, 8, None, False, None),

So DeepSeek-V3: $d = 7168$, $d_{ff,\text{moe}} = 2048$, $E = 256$, $k = 8$. Qwen3-30B-A3B: $d = 2048$, $d_{ff,\text{moe}} = 768$, $E = 128$, $k = 8$. Layer counts (61 with the first 3 dense for V3, 48 for Qwen3) and the group-limited routing parameters ($n_{\text{group}} = 8$, $\text{topk\_group} = 4$) come from the published configs that §4.4 read; SGLang's config loader confirms the V3 template's topk_group=4 in passing (python/sglang/srt/configs/model_config.py:L411-L420) and both engines' routing tests instantiate num_expert_group=8, topk_group=4 (tests/kernels/moe/test_routing.py:L83-L84 and test/registered/moe/test_topk_renormalize_degenerate.py:L102-L103).

An expert is three matrices — gate, up, down — stored as a fused w13_weight of $[E,\,2 d_{ff},\, d]$ plus a w2_weight of $[E,\, d,\, d_{ff}]$:

        if self.moe.is_act_and_mul:
            w13_up_dim = 2 * intermediate_size_per_partition
        else:
            w13_up_dim = intermediate_size_per_partition
        # Fused gate_up_proj (column parallel)
        w13_weight = torch.nn.Parameter(
            torch.empty(
                num_experts,
                w13_up_dim,
                hidden_size,
                dtype=params_dtype,
            ),
            requires_grad=False,
        )
# ...
        # down_proj (row parallel)
        w2_weight = torch.nn.Parameter(
            torch.empty(
                num_experts,
                hidden_size,
                intermediate_size_per_partition,
                dtype=params_dtype,
            ),

So one expert holds $3 d\, d_{ff}$ parameters and costs $6 d\, d_{ff}$ FLOPs per token that routes to it. For DeepSeek-V3 that is $3 \times 7168 \times 2048 = 44.04$ M parameters and $6 \times 7168 \times 2048 = 88.08$ MFLOP per token-expert.

Two parameter counts, two bills

The three derived rows for V3 sum to 34.40 B against a published 37 B activated; the balance is the three dense MLP layers and the embedding and output matrices, which I did not count. The check matters — a parameter count that does not reconcile with the model card means you read a shape wrong.

The FLOP bill follows the active count. A decode step does $F \approx 2 N_{\text{active}}$ FLOPs per token: 74 GFLOP for DeepSeek-V3, 6.6 GFLOP for Qwen3-30B-A3B — the profiles of a 37 B and a 3 B dense model. The HBM bill follows neither. It follows how many distinct experts the batch touched.

How many experts does a batch touch?

Assume uniform routing and independent tokens. This is a reference model, not a conservative bound on all costs: skew may reduce distinct weights read while increasing the busiest expert's latency. Under this assumption an expert is missed by a token with probability 1 - k/E, giving:

This is the same coupon-collector shape as $\rho$ in §5.3, applied to experts instead of ranks. Its two limits are the story: $E_{\text{touched}}(1) = k$, and $E_{\text{touched}}(B) \to E$ fast. The half-way batch — where you already stream half the layer — is $B_{1/2} = \ln(1/2)\,/\,\ln(1 - k/E)$: 21.8 tokens for DeepSeek-V3, 10.7 for Qwen3-30B-A3B. You reach 95% of experts at $B = 94$ and $B = 46$ respectively.

Arithmetic intensity versus batch

Now put both bills together for one MoE layer, at batch $B$, with $b_w$ bytes per weight element. FLOPs are $F = B\,k\,6\,d\,d_{ff}$. Compulsory weight traffic is $Q = E_{\text{touched}}(B)\cdot 3\,d\,d_{ff}\,b_w$; activations are $O(Bd)$ and negligible against $O(E d\, d_{ff})$. Dividing, $d$ and $d_{ff}$ cancel completely:

Two limits, and both are load-bearing. At $B = 1$, $E_{\text{touched}} = k$ and $I = 2/b_w$ — exactly the dense decode GEMV intensity from §0.4. At batch 1 an MoE layer and a dense layer are the same kind of object: you read a weight byte and do one multiply-add with it. As $B \to \infty$, $E_{\text{touched}} \to E$ and

The asymptotic slope is the dense slope divided by $E/k$. For DeepSeek-V3 that ratio is 32; for Qwen3-30B-A3B, 16. Setting $I_{\text{MoE}} = I^{*}$ and solving gives the batch at which the expert GEMM finally becomes compute-bound:

Which is §4.4's $B^{*}_{\text{MoE}} = M^{*} \cdot E/k$ arrived at from the other end: the same 296-token dense crossover, multiplied by the same $E/k = 32$.

What that costs in wall-clock, on one H100

DeepSeek-V3 in fp8 is 671 GB and does not fit on one GPU, so take Qwen3-30B-A3B, which at 28.99 GB of fp8 expert weights does. Per decode step, expert-weight streaming alone, at the H100's 3.35 TB/s:

Read the last two columns together. Step time stops growing at batch 128 — past that point every additional token is free in bandwidth terms, which is why MoE serving wants the largest batch it can assemble. The batch-1 row is the other half: 541 µs of expert weights for a 30 B model, against the book's dense-8 B bf16 decode floor of 4.48 ms. At batch 1 a sparse 30 B model is cheaper than a dense 8 B one; at batch 512 it costs the same as a dense 30 B one. MoE has the small model's compute profile always, the small model's memory profile at batch 1, and the big model's memory profile everywhere else.

The gate is a single linear layer, $d \times E$, with no bias — 1.84 M parameters for DeepSeek-V3, 0.02% of the layer. Everything interesting is what happens to its output. vLLM's generic path is fused_topk, which delegates to a CUDA kernel and branches only on the scoring function:

def fused_topk(
    hidden_states: torch.Tensor,
    gating_output: torch.Tensor,
    topk: int,
    renormalize: bool,
    indices_type: torch.dtype | None = None,
    scoring_func: str = "softmax",
) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
    assert hidden_states.size(0) == gating_output.size(0), "Number of tokens mismatch"

    M, _ = hidden_states.size()

    topk_weights = torch.empty(
        M, topk, dtype=torch.float32, device=hidden_states.device
    )
    topk_ids = torch.empty(
        M,
        topk,
        dtype=torch.int32 if indices_type is None else indices_type,
        device=hidden_states.device,
    )

Note the two live choices. Softmax before or after top-k: Mixtral-style routing softmaxes the full logit vector first, so the $k$ selected weights do not sum to 1; renormalize=True divides by their sum afterwards, which is equivalent to softmaxing only over the selected $k$. DeepSeek-V3 instead scores with sigmoid — independent per-expert gates, not a distribution — and then renormalises. The model config carries the choice (scoring_func, norm_topk_prob) and vLLM passes it straight through (vllm/model_executor/models/deepseek_v2.py:L370-L386).

DeepSeek's actual routing is the grouped, bias-corrected variant, and this is the code worth reading carefully:

    if scoring_func == "softmax":
        scores = torch.softmax(gating_output, dim=-1)
    elif scoring_func == "sigmoid":
        scores = gating_output.sigmoid()
    else:
        raise ValueError(f"Unsupported scoring function: {scoring_func}")

    num_token = scores.size(0)
    if e_score_correction_bias is not None:
        # Store original scores before applying correction bias. We use biased
        # scores for expert selection but original scores for routing weights
        original_scores = scores
        scores = scores + e_score_correction_bias.unsqueeze(0)
        group_scores = (
            scores.view(num_token, num_expert_group, -1).topk(2, dim=-1)[0].sum(dim=-1)
        )
# ...
    group_idx = torch.topk(group_scores, k=topk_group, dim=-1, sorted=use_sorted)[
        1
    ]  # [n, top_k_group]
    group_mask = torch.zeros_like(group_scores)  # [n, n_group]
    group_mask.scatter_(1, group_idx, 1)  # [n, n_group]
    score_mask = (
        group_mask.unsqueeze(-1)
        .expand(num_token, num_expert_group, scores.size(-1) // num_expert_group)
        .reshape(num_token, -1)
    )  # [n, e]
    tmp_scores = scores.masked_fill(~score_mask.bool(), float("-inf"))  # [n, e]

    if e_score_correction_bias is not None:
        topk_ids = torch.topk(tmp_scores, k=topk, dim=-1, sorted=use_sorted)[1]
        # Use original unbiased scores for the routing weights
        topk_weights = original_scores.gather(1, topk_ids)
    else:
        topk_weights, topk_ids = torch.topk(
            tmp_scores, k=topk, dim=-1, sorted=use_sorted
        )

    if renormalize:
        topk_weights = topk_weights / topk_weights.sum(dim=-1, keepdim=True)

Three mechanisms, stacked:

SGLang's reference implementation is the same algorithm with one extra move — it can splice the shared expert in as a fake routed expert:

    scores = gating_output.sigmoid()
    num_token = scores.shape[0]
    num_experts = scores.shape[1]
    scores_for_choice = scores.view(num_token, -1) + correction_bias.unsqueeze(0)
    group_scores = (
        scores_for_choice.view(num_token, num_expert_group, -1)
        .topk(2, dim=-1)[0]
        .sum(dim=-1)
    )  # [n, n_group]
# ...
    topk_weights = scores.gather(1, topk_ids)

    if num_fused_shared_experts:
        topk_ids[:, -1] = torch.randint(
            low=num_experts,
            high=num_experts + num_fused_shared_experts,
            size=(topk_ids.size(0),),
            dtype=topk_ids.dtype,
            device=topk_ids.device,
        )
        if routed_scaling_factor is not None:
            topk_weights[:, -1] = (
                topk_weights[:, :-1].sum(dim=-1) / routed_scaling_factor
            )

That is worth understanding before §7: the shared expert is appended as expert id 256 and given a synthetic weight, so top-8 becomes top-9 and the shared FFN disappears into the same grouped GEMM instead of being a separate launch.

After routing, expert $e$ owns $m_e$ rows, data-dependent, with $\sum_e m_e = Bk$. A GEMM wants one $M$. Padding every expert to the batch size is out — for DeepSeek-V3 at $B = 512$ that is $256 \times 512 = 131{,}072$ rows to do 4,096 rows of work. Looping over experts in Python is also out: 256 launches per layer over 58 layers is 14,848 kernel launches per decode step, against a step budget of a few milliseconds. Two workable answers remain — index into the activations in place, or physically permute them into per-expert segments. Both engines implement both; which one you get is a dispatch decision, and the defaults below are the ones that matter.

vLLM's Triton path: align, then index

vLLM's moe_align_block_size is the clearest single artefact in either codebase for understanding MoE kernels, because its docstring works the example:

    Example:
    Given topk_ids = [[2, 3, 4], [1, 2, 4], [1, 3, 4], [1, 2, 3]],
    block_size = 4, and num_experts = 4:
    - We initially have 12 tokens (after repeating 'top_k' times) and 4 experts,
        with each expert needing to process 3 tokens.
    - As block_size is 4, we pad 1 token for each expert.
    - First, flatten topk_ids to [2, 3, 4, 1, 2, 4, 1, 3, 4, 1, 2, 3].
    - Then append padding tokens [12, 12, 12, 12] for each block.
    - After sorting by expert index, we obtain token_ids
        [3, 6, 9, 12, 0, 4, 10, 12, 1, 7, 11, 12, 2, 5, 8, 12].
        Tokens 12 are non-existent (padding) and are ignored in
        the subsequent matrix multiplication.
    - The padding ensures that the total number of tokens is now divisible
        by block_size for proper block matrix operations.

The upstream example uses labels 1 through 4. Executable zero-based IDs for four experts must be 0 through 3: subtract one from each expert ID; assignment indices are unchanged.

SGLang carries a byte-identical docstring at python/sglang/srt/layers/moe/moe_runner/triton_utils/moe_align_block_size.py:L60-L73 — the Triton MoE path has a shared ancestor and the two projects have been trading patches on it since. The CUDA implementation counts per-expert hits in shared memory, ceils each count to the block size, block-scans them into offsets, and then has every thread atomicAdd its way into the right slot:

  int expert_count = 0;
  int expert_id = threadIdx.x;
  if (expert_id < num_experts) {
    int warp_idx = expert_id / experts_per_warp;
    int expert_offset = expert_id % experts_per_warp;
    expert_count = shared_counts[warp_idx * experts_per_warp + expert_offset];
    expert_count = CEILDIV(expert_count, block_size) * block_size;
  }

  int cumsum_val;
  BlockScan(temp_storage).ExclusiveSum(expert_count, cumsum_val);
  if (expert_id <= num_experts) {
    cumsum[cumsum_offset + expert_id] = cumsum_val;
  }
# ...
  if (threadIdx.x < num_experts) {
    for (int i = cumsum[cumsum_offset + threadIdx.x];
         i < cumsum[cumsum_offset + threadIdx.x + 1]; i += block_size) {
      expert_ids[expert_ids_offset + i / block_size] = threadIdx.x;
    }
  }

  // Fill remaining expert_ids with -1
  const size_t fill_start_idx =
      cumsum[cumsum_offset + num_experts] / block_size + threadIdx.x;
  for (size_t i = fill_start_idx; i < max_num_m_blocks; i += blockDim.x) {
    expert_ids[expert_ids_offset + i] = inactive_expert_id;
  }

The line that matters is expert_count = CEILDIV(expert_count, block_size) * block_size. Each expert's segment is rounded up to a whole tile. That single ceiling is the source of most of MoE's tensor-core waste at decode, quantified below.

The payoff is that the activations are never physically permuted. The Triton kernel reads the sorted index list and divides by top_k to recover the original row:

    offs = tl.arange(0, BLOCK_SIZE_M).to(tl.int64)
    num_tokens_post_padded = tl.load(num_tokens_post_padded_ptr)
    if pid_m * BLOCK_SIZE_M >= num_tokens_post_padded:
        return
    if not naive_block_assignment:
        offs_token_id = pid_m * BLOCK_SIZE_M + offs
        offs_token = tl.load(sorted_token_ids_ptr + offs_token_id)
# ...
    token_mask = offs_token < num_valid_tokens

    off_experts = tl.load(expert_ids_ptr + pid_m).to(tl.int64)
# ...
        a_ptrs = a_ptr + (
            offs_token[:, None] // top_k * stride_am + offs_k[None, :] * stride_ak
        )
        b_ptrs = (
            b_ptr
            + off_experts * stride_be
            + (offs_bn[:, None] * stride_bn + offs_k[None, :] * stride_bk)
        )

Each M-block reads one expert id, offsets into B by off_experts * stride_be, and gathers its $\le$ BLOCK_SIZE_M activation rows through offs_token // top_k. Invalid slots are masked. The unpermute is equally cheap: the output cache is $[M, \text{topk}, N]$, and ops.moe_sum reduces the topk axis (vllm/model_executor/layers/fused_moe/fused_moe.py:L1854-L1857) — no gather, no scatter, just a strided add. The routing weight is not applied here: the second grouped GEMM is dispatched with mul_routed_weight = not apply_router_weight_on_input, so w2's kernel scales each row by its topk_weight on the way out and moe_sum is a plain unweighted sum (:L1830-L1857).

At very small batches vLLM skips the alignment kernel entirely:

    """Prepare expert assignments for the aligned and low-latency Triton paths."""
    # SPARSITY_FACTOR is a heuristic margin ensuring tokens_in_chunk * top_k
    # activates only a small fraction of total experts
    # Skips moe_align_block_size and activates the `sorted_token_ids is None`
    # path of the fused_moe_kernel kernel
    naive_block_assignment = (
        expert_map is None
        and num_tokens * top_k_num * 4 <= global_num_experts
        and not (
            (use_int8_w8a16 or use_int4_w4a16)
            and block_shape is not None
            and block_shape[1] > 0
        )
    )

For DeepSeek-V3, $B \cdot 8 \cdot 4 \le 256$ means $B \le 8$: at batch 8 or below, every M-block is one token and the sort is pure overhead. Note the first conjunct, expert_map is None — under expert parallelism there is a map to apply, so the shortcut is off no matter how small the batch, and the alignment kernel runs at every batch size.

SGLang under DeepGEMM: permute for real, because DeepGEMM demands it

When SGLang's DeepGEMM runner is selected — not the default, per the callout above — the contiguous grouped GEMM requires each expert's segment to start at a 128-row boundary in one packed buffer. So it physically copies:

    # The compact layout avoids scaling masked buffers with the expert count.
    # Scatter and post-permute skip non-local experts mapped to -1.
    block_e = 128
    num_experts = runner_config.num_local_experts
    num_assignments = topk_ids.numel()
    all_tokens = _get_compact_all_tokens(num_assignments, num_experts, block_e)

    tokens_per_expert, unused_masked_dst = fused_moe_dispatch_index(
        topk_ids, num_experts, 1
    )
    dispose_tensor(unused_masked_dst)
    valid_tokens_per_expert = tokens_per_expert
    tokens_per_expert = (ceil_div(tokens_per_expert, block_e) * block_e).to(torch.int32)
    # Keep graph-static shapes by appending padding to the final segment.
    # Its m_indices stay -1, so DeepGEMM skips those rows.
    tokens_per_expert[-1].add_(all_tokens - tokens_per_expert.sum())

Then ep_scatter writes the permuted rows and the segment map, the grouped GEMM runs once over m_indices, and post_reorder_deepgemm reverses it:

    expert_start_loc = torch.empty(
        num_experts, device=hidden_states_device, dtype=torch.int32
    )
    m_indices = torch.empty(all_tokens, device=hidden_states_device, dtype=torch.int32)
    src2dst = torch.empty_like(topk_ids, dtype=torch.int32)
    ep_scatter(
        packed_input_source,
        packed_input_source_scale,
        topk_ids,
        tokens_per_expert,
        valid_tokens_per_expert,
        expert_start_loc,
        packed_input,
        packed_input_scale,
        m_indices,
        src2dst,
# ...
    post_reorder_deepgemm(
        runner_output.hidden_states,
        output,
        src2dst,
        topk_ids,
        topk_weights,
        runner_config.top_k,
        hidden_states_shape[0],
        hidden_states_shape[1],

The tradeoff is explicit. vLLM's index-gather avoids a full copy of the activations — for DeepSeek-V3 at $B = 512$ that copy is $Bk \times d = 29.4$ MB of fp8 written and read again. SGLang pays it to get contiguous, aligned inputs, which is what lets it hand the GEMM to DeepGEMM's hand-tuned Hopper kernels instead of Triton. Whether that wins on your shapes is a measurement, and neither tree contains one.

The padding tax, quantified

Define the padding factor as padded rows over real rows. When each expert's mean row count $\bar m = Bk/E$ is well below the tile height $M_{\text{blk}}$, every touched expert costs a full tile:

This is a tensor-core tax, not a bandwidth tax — the padded rows are zeros that still occupy MMA slots. And it explains an artefact you can read straight out of vLLM's repository. The project ships autotuned Triton configs per $(E, N, \text{device})$; here is the tile height it found for DeepSeek-V3 at TP=4 on an H100:

The autotuner independently rediscovered §4.4's $M_e = Bk/E$: from batch 256 upward the winning tile is consistently $2 M_e$, and below batch 256 it floors at 16 because Triton cannot go lower usefully. A dense GEMM at batch 2048 would want a 128-row tile; the MoE layer at batch 2048 also wants 128 — because its effective $M$ is 64, not 2048. The heuristic fallback encodes the same insight in a comment: "Use a small M tile for decode-like batches where tokens are spread thin across experts" (vllm/model_executor/layers/fused_moe/fused_moe.py:L1320-L1324).

The training literature will tell you that MoE layers have a capacity factor: each expert gets a fixed buffer of $C = \lceil \alpha B k / E \rceil$ rows, and tokens that arrive at a full expert are dropped — their contribution from that expert is zero and only the residual carries them forward. GShard and Switch Transformer both work that way, and it exists because training wants static shapes.

That decision has a real cost, and you can watch it being paid. Three buffer-sizing strategies appear in the two trees. Contiguous/compact sizes to what actually happened: vLLM allocates max_num_tokens_padded = topk_ids.numel() + num_experts * (block_size - 1) (vllm/model_executor/layers/fused_moe/moe_align_block_size.py:L74), so the worst case is one wasted partial tile per expert — bounded and small. Batched/masked gives every expert a rectangle of $m_{\max}$ rows; vLLM's BatchedPrepareAndFinalize, the reference format for all-to-all backends, allocates $[E_{\text{local}},\ m_{\max},\ d]$ with $m_{\max}$ set to the scheduler's entire token budget (vllm/model_executor/layers/fused_moe/layer.py:L259, L343) — a capacity factor of $E/k$ in effect, enough that all $B$ tokens could pile onto one expert. Correct by construction, and expensive. Measured cap is SGLang's fix after hitting exactly that wall — opt-in: SGLANG_OPT_DG_MASKED_M_CAP defaults to False (python/sglang/srt/environ.py:L989), so on a stock launch the uncapped $m_{\max}$ is what allocates. Its comment is the most instructive text in either codebase on this topic:

    # For masked grouped GEMM, shape M should be multiple of the block M (current block M: {block_m}) https://github.com/deepseek-ai/DeepGEMM/blob/main/deep_gemm/jit_kernels/m_grouped_gemm.py#L165
    m_max = (hidden_states.size(0) // 256 + 1) * 256
    if (
        envs.SGLANG_OPT_DG_MASKED_M_CAP.get()
        and not torch.cuda.is_current_stream_capturing()
    ):
# ...
        # m_max reserves capacity for ALL rank tokens in EVERY local expert:
        # the [num_local_experts, m_max, *] masked-GEMM intermediates reach
        # 7+ GiB per 32k-token chunk and OOM saturated serving.  The hottest
        # expert only ever holds max(masked_m) rows, so cap the padded
        # capacity there (rounded up to the DeepGEMM block-M).  Costs one
        # probe dispatch-index launch + one D2H sync per MoE layer;
        # correctness is unconditional (m_cap >= max(masked_m) by
        # construction, and the final src2dst below is built with the same
        # capped stride).
        masked_m_probe, _ = fused_moe_dispatch_index(topk_ids, num_local_experts, m_max)
        m_cap = (int(masked_m_probe.max().item()) + 255) // 256 * 256
        m_max = min(m_max, max(m_cap, 256))
    expected_m = (topk_ids.numel() - 1) // num_local_experts + 1

    masked_m, src2dst = fused_moe_dispatch_index(topk_ids, num_local_experts, m_max)

    gateup_input = torch.empty(
        (num_local_experts, m_max, hidden_states.size(1)),
        device=hidden_states.device,
        dtype=output_dtype,
    )

Read the tradeoff: a device-to-host sync per MoE layer, deliberately accepted, to avoid allocating for a worst case that never happens — and two guards. The env flag is off unless you set it, and even when set the sync is illegal during CUDA-graph capture, so decode keeps the uncapped $m_{\max}$. SGLang also picks between the masked and compact layouts by estimating each one's peak bytes against a memory budget (python/sglang/srt/layers/moe/moe_runner/deep_gemm.py:L117-L152, landed in SGLang #33474).

The cited serving paths preserve all routed token assignments rather than using training-style token dropping. They still need static allocations during graph capture: worst-case/reserved buffers, masks and padding coexist with dynamic logical row counts. Overflow must trigger resizing, fallback or explicit failure, not silent token loss.

DeepSeek-V3 has one always-on expert per MoE layer alongside the 256 routed ones; Qwen3-30B-A3B has none (shared_expert_intermediate_size defaults to 0 and the module is skipped, vllm/model_executor/models/qwen3_moe.py:L175-L197). The argument is redundancy: without one, every routed expert must independently learn the general-purpose transformations all tokens need, spending capacity on $E$ copies of the same thing.

At batch 1, one same-size shared expert adds 1/8 = 12.5% to the eight routed experts' weight bytes, and accounts for 1/9 = 11.1% of their combined total. When all 256 routed experts are touched, it adds 1/256 and contributes 1/257 of the total. These byte fractions do not prove negligible latency: shared-expert compute, launches and overlap still matter. Placement and fusion decide whether work is local and whether a separate launch is needed. Both engines can fuse it into the routed path rather than launching it separately, but only one of them does so on NVIDIA by default. SGLang remaps the checkpoint tensor into an extra expert slot:

        # num_fused_shared_experts drives weight remapping in deepseek_weight_loader:
        # mlp.shared_experts → mlp.experts.256 when > 0.
        self.num_fused_shared_experts = 0 if _fusion_disabled else n_shared_experts

        # DeepEP and MegaMOE shared expert fusion: shared expert is fused into
        # the same MoE kernel as a local expert at each EP rank. Expert layout
        # is expanded from 256 routed to 256+EP_size (e.g. 272 for EP=16).
        _uses_per_rank_shared_slots = has_per_rank_fused_shared_slots(
            self.num_fused_shared_experts
        )

        if _uses_per_rank_shared_slots:
            # 256 routed + EP_size shared slots = 272 experts total (for EP=16)
            num_experts_for_moe = config.n_routed_experts + self.moe_ep_size
            top_k_for_moe = config.num_experts_per_tok + 1  # 8 routed + 1 shared

vLLM has the same lever — fuse_shared_experts on the MoE factory (vllm/model_executor/models/deepseek_v2.py:L391-L394) — but at this SHA it is never pulled on NVIDIA. resolve_layer_fused_shared_expert gates entirely on rocm_aiter_ops.is_fusion_moe_shared_experts_enabled() (vllm/model_executor/layers/fused_moe/utils.py:L73-L85), and that classmethod is wrapped in @if_aiter_supported, which returns None off ROCm gfx9 (vllm/_aiter_ops.py:L195-L207, L1836-L1838). So on an H100 vLLM builds a plain DeepseekV2MLP at $\text{moe\_intermediate\_size} \times n_{\text{shared}}$ (vllm/model_executor/models/deepseek_v2.py:L355-L368) and hands it to the MoE layer as a separate module.

Expert weights dominate everything. DeepSeek-V3's 653.9 GB of routed experts in fp8 is 97.5% of the model; in bf16 it is 1,308 GB. Put that against §2.1's budget. Mind the units, because this table is in decimal GB and the card is not: an “80 GB” H100 is 79.65 GiB. At gpu_memory_utilization=0.92 that is 73.28 GiB, and after the 6.0 GiB §2.1 charges for activations, graphs and non-torch allocations, 67.28 GiB is left for weights plus KV — which, in the decimal GB this table uses, is ~72 GB per GPU. (That it also equals a careless 80 × 0.9 is a coincidence of two unit errors cancelling; the derivation above is the real one.)

Quantization on a dense model buys latency in the memory-bound regime and stops mattering once you are past the ridge. On an MoE model it buys GPU count, and — per the derivation in §3 — the expert layers do not reach the ridge until batch 9,456, so the latency win never runs out either. That is why the natural pairing exists, and why both engines have quantization schemes that touch only the MoE. vLLM's moe_wna16 is one: its get_quant_method routes RoutedExperts to MoeWNA16Method while sending every LinearBase off to the underlying GPTQ or AWQ config (vllm/model_executor/layers/quantization/moe_wna16.py:L171-L197). §4.4 owns the decision procedure and the accuracy side; this chapter only supplies the reason it works.

§5.3 treats imbalance as a placement problem: which rank holds which expert, measured as balancedness, fixed by EPLB. On one device there is no placement to fix, and imbalance still costs you — through the grouped GEMM's geometry.

Two distinct mechanisms:

The masked column is the one to internalise: the slowest group sets the time, exactly as it does across ranks in §5.3, but here on a single device with no placement lever to pull. Your only moves are the layout choice and the batch composition. A perfectly balanced batch at $B = 4096$ pays 1.0× in the contiguous layout and 1.0× in the masked one; the batch drawn above pays 1.42× and 2.00×.

DeepSeek-V3, vLLM, single GPU (imagine it fits), fp8 block-quantised weights. One decode step with 64 sequences enters one MoE layer.

1. DeepseekV2MoE.forward reshapes to [64, 7168] and calls the experts module with router_logits=hidden_states — the gate GEMM happens inside (vllm/model_executor/models/deepseek_v2.py:L411-L421).
2. grouped_topk runs: sigmoid on the [64, 256] logits, add e_score_correction_bias, score the 8 groups by their top-2 sum, keep the best 4, mask the rest to $-\infty$, take top-8, gather weights from the unbiased scores, renormalise, scale (vllm/model_executor/layers/fused_moe/router/grouped_topk_router.py:L112-L161). Output topk_ids [64, 8], topk_weights [64, 8].
3. fused_experts_impl quantises activations to fp8, then calls _prepare_expert_assignment. The naive shortcut needs $64 \times 8 \times 4 \le 256$, which is false, so it falls through to moe_align_block_size (:L1557-L1584).
4. The CUDA kernel builds the layout. With BLOCK_SIZE_M = 16, max_num_tokens_padded = 512 + 256 × 15 = 4{,}352 and expert_ids is 272 entries; the 512 assignments land on ~222 distinct experts, so num_tokens_post_padded comes out near 3,552 — 6.9 padded rows per real row (derived). Absent experts get expert_ids = -1.
5. dispatch_fused_moe_kernel launches over $\lceil 4352/16 \rceil \times \lceil 4096/64 \rceil$ program ids. Each loads its expert id; if it is $-1$ it calls write_zeros_to_output and returns (:L423-L440). Otherwise it gathers up to 16 rows via offs_token // top_k and streams w13[e]. Across the layer, ~222 experts × 44.04 MB = 9.78 GB streamed for 64 tokens.
6. apply_moe_activation runs SiLU-and-mul, halving 4096 to 2048; the result is re-quantised and fed to a second dispatch_fused_moe_kernel against w2 with top_k=1, so that gather is the identity, and with mul_routed_weight on, so this kernel is where topk_weights is applied (:L1815-L1852).
7. ops.moe_sum then adds the already-scaled [64, 8, 7168] down to [64, 7168] (:L1854-L1857); the shared expert and residual are added. Net for the layer's routed experts: $k \cdot 2 \cdot 44.04\,\text{M} = 705$ MFLOP per token against 9.78 GB of weight reads for the 64-token step — an intensity of 4.6 FLOP/byte on a machine whose fp8 ridge is 591.

Selection bias may leave the raw selected scores unchanged while changing the selected expert and therefore the output. For top-1 normalized routing the selected weight becomes one. The example also groups assignments by expert, scatters their weighted outputs back, and compares with a direct per-token reference. Empty experts and allocation capacity do not contribute outputs.

import numpy as np

x = np.array([[1., 2.], [3., 4.], [-1., 1.]])
weights = np.array([np.eye(2), 2*np.eye(2), -np.eye(2)])
ids = np.array([[0, 1], [1, 2], [0, 2]])
routing = np.array([[0.7, 0.3], [0.4, 0.6], [0.5, 0.5]])
direct = sum(routing[:, j, None] *
    np.stack([x[t] @ weights[ids[t, j]].T for t in range(len(x))])
    for j in range(2))
grouped = np.zeros_like(x)
for expert in range(len(weights)):
    tokens, slots = np.where(ids == expert)
    np.add.at(grouped, tokens,
              routing[tokens, slots, None] * (x[tokens] @ weights[expert].T))
np.testing.assert_allclose(grouped, direct)
scores = np.array([0.6, 0.4])
before = np.argmax(scores)
after = np.argmax(scores + np.array([0., 0.3]))
assert before != after
assert not np.allclose(x[0] @ weights[before].T, x[0] @ weights[after].T)
print("Grouped scatter matches direct outputs; selection bias changes the function.")

Everything below runs without a GPU except the last item.

V=~/Documents/other_git_repos/vllm

# 1. The expert-touch curve and the intensity curve, for any (E, k).
python3 - <<'PY'
E, k, bw, ridge = 256, 8, 1, 591          # DeepSeek-V3, fp8, H100
for B in (1, 8, 32, 64, 128, 256, 512, 4096):
    touched = E * (1 - (1 - k / E) ** B)
    print(f"B={B:5d}  touched={touched:7.1f}  I={2*B*k/(bw*touched):8.2f}")
print("ridge at B =", bw * ridge * E / (2 * k))
PY

# 2. Read the autotuner's answer for the tile height.
python3 -c "
import json, pathlib
p = pathlib.Path('$V/vllm/model_executor/layers/fused_moe/configs')
d = json.load(open(p / 'E=256,N=512,device_name=NVIDIA_H100_80GB_HBM3.json'))
for b, c in d.items():
    print(f\"B={b:>5}  BLOCK_SIZE_M={c['BLOCK_SIZE_M']:>4}  M_e={int(b)*8/256:7.1f}\")
"

# 3. Confirm the no-token-dropping claim yourself.
grep -rn "capacity_factor\|drop_tokens" $V/vllm/ | grep -v vision | grep -v arctic

On a GPU, the flag to flip is the MoE backend. For block-fp8 weights on a single-node Hopper box both engines default to the index-gather Triton path, so you have to ask for the other one. On vLLM, --moe-backend deep_gemm selects it directly, while --enable-expert-parallel (default False, vllm/config/parallel.py:L165, and it needs more than one GPU) instead promotes FlashInfer CUTLASS on Hopper, which needs flashinfer installed or it falls through. On SGLang, --moe-runner-backend deep_gemm is the direct route and --moe-a2a-backend deepep (default none) is the indirect one — the latter additionally needs the external deep_ep package and multiple ranks, and both need deep_gemm importable. Measure with the standard benchmark harness at fixed input and output lengths, sweeping --max-num-seqs and measure actual resident batch sizes. The reference byte model predicts amortized cost falling as occupancy grows. Once distinct weight bytes saturate, constant step time would mean per-token cost continues falling as 1/B, until compute or other work grows. A plateau in per-token cost is not implied by saturated weight traffic. §10.1 owns methodology.

• Shazeer et al., Outrageously Large Neural Networks: The Sparsely-Gated Mixture-of-Experts Layer (2017) — the top-k gate and the load-balancing loss, in their original form.
• GShard and Switch Transformer — where the capacity factor and token dropping come from. Read them knowing that §6 above shows inference engines discarded both.
• MegaBlocks: Efficient Sparse Training with Mixture-of-Experts — the block-sparse reformulation that made dropless MoE practical, and the direct ancestor of the alignment-plus-grouped-GEMM structure both engines use.
• Wang et al., Auxiliary-Loss-Free Load Balancing Strategy for Mixture-of-Experts (2024) — the e_score_correction_bias mechanism, and why separating selection scores from routing weights matters.
• DeepSeek-V3 Technical Report — 671 B total, 37 B activated, 256 routed experts plus 1 shared, group-limited routing. The model card carries the config this chapter's shapes come from; vLLM's benchmark table is the in-tree copy.
• Qwen3-30B-A3B model card — 128 experts, top-8, no shared expert, every layer sparse.
• DeepGEMM — the contiguous and masked grouped-GEMM layouts SGLang targets. SGLang's own source links deep_gemm/jit_kernels/m_grouped_gemm.py#L165 for the block-M alignment rule.
• PRs worth reading: vLLM #47785 (topk_ids padding in the align-sum kernel) and #46643 (vectorised fp32 moe_sum, arbitrary top-k); SGLang #33128 (DeepGEMM for standard MoE dispatch) and #33474 (choosing the layout by memory budget).
• In this book: §5.3 for expert parallelism, the all-to-all and EPLB; §4.4 for the quantization decision on expert weights; §7.2 for MLA, the other half of what makes DeepSeek-V3 servable.