SGLang 源码完全拆解（二十）：门控 Router 与按专家对齐——三份路由账本、五种输出格式与 moe_align_block_size 的排序账

承接：从”专家放在哪”回到”token 去找谁”

SGLang 子系列走到第十九篇为止，请求主线与并行主线都已经闭合：请求经 TokenizerManager 分词、由 ZMQ 送进 Scheduler，调度器把 token 预算与 KV 页交给 PrefillAdder 与 RadixCache，TpModelWorker 与 ModelRunner 组装 ForwardBatch 交给各层执行；注意力侧的后端、KV 池、前缀缓存都有独立篇目；并行侧先讲 DP Attention 与专家并行（SGLang 源码完全拆解（十七）：DP Attention 与专家并行），再讲弹性 EP 与故障恢复（SGLang 源码完全拆解（十八）：弹性 EP 与故障恢复）。

这两篇解决的是 MoE 的物理侧：专家由谁持有、放在哪张卡上、重排时怎么搬。但一个 token 进入 MoE 层后，第一件事是回答逻辑侧的问题：这一层的 256（或 512、896）个专家里，它应该去投哪几个。SGLang 里这个决定主要由 python/sglang/srt/layers/moe/topk.py（约 3200 行）做出，打分与 top-k 的具体内核则落在 python/sglang/kernels/ops/moe/ 下的 JIT kernel 里。

本篇是第一篇专讲 MoE 的单框架文章，聚焦路由的前半程：从门控 logit 到”可以交给分桶内核”的 topk_ids 与 topk_weights，再往后一小步到 moe_align_block_size 产出的按专家排序表。分组 GEMM 本体、EP 通信、EPLB 重排作为下一阶段的主题；这里只到”账本交出去”为止。

之所以值得专门写这一层，是因为它看起来只是”取前 k”，实际却把六件容易混在一起的事压在几行调用里：

1. 门控分数怎么从 logit 变成”愿望值”（softmax / sigmoid / sqrtsoftplus）；
2. DeepSeek 式的分组选组、修正偏置”只参与选、不参与算”；
3. 选定后要不要归一、归约的起点与 epsilon；
4. 共享专家是拼进 topk 输出、还是另外走一条路；
5. 逻辑专家号什么时候变成物理槽位（EPLB）；
6. 交给下游时，是标准三元组、紧凑压缩格式、还是干脆不路由。

下面按”打分 → 选择 → 后处理 → 对齐”四段走一遍源码，所有代码引自仓库当前版本（python/sglang/ 目录，2026 年 10 月的开发分支），少量简化只体现在省略号，不改变语义。

图中的红色虚线是整篇最需要记住的一条路径：修正偏置进入”选组 / 选专家”的比较，却在”权重回取”时从 activated（未修正分数）取值；下游的 topk_weights 因此只反映模型原始分数，偏置不会改变输出幅度。

门控这一步在推理里要满足什么约束

先把”为什么不能简单写torch.topk”说清楚，后文所有取舍都围绕这几点。

稠密 FFN 每层是两个矩阵；MoE 把它换成 (N) 个专家加一张门控矩阵 (W_g \in \mathbb{R}^{N \times d})。一个 token 的隐状态 (x \in \mathbb{R}^d)：

(g) 的第 (j) 个分量表示”这个 token 有多想去专家 (j)“。从 (g) 到真正算完，中间有选择（256 选 8）、权重（归一）、分桶、逐专家 GEMM、加权合并几步。门控层只负责前两步，但它的输出格式直接决定后面能不能按块切。

Mixtral 的 8 选 2 是最朴素的路由：一个 token 一个 router 向量，softmax 后取前 2 名，两个专家各自跑一遍 FFN，再按权重相加。SGLang 的 fused_topk 负责的正是图中”router → 2 个专家 → 加权”里的前后两段，中间那段 GEMM 由后续分组内核接管。

推理与训练的不同在形状：训练的 batch 大且固定，可以容忍每一步都做 all-to-all；推理时 batch 可能是 1 条 prefill 的两万 token，也可能是 32 条 decode 的 32 个 token，每层路由出来的专家分布都不一样。由此得到三条约束：

• 不能在主机上等 topk 结果。一旦把 topk_ids 拷回 CPU 来决定下一步，解码每 step 就多一次 kernel 加一次同步，而 decode 一次 step 也许只有几百微秒；
• 输出必须是下游内核能直接下标访问的布局，不能”先产出分数、再让下游自己找”；
• 同一请求在任意 batch 下结果尽量稳定。路由是排序，平局怎么破、归约顺序怎么定，都会让同一 token 在不同 batch 大小下走出不同的专家。

topk.py 的结构就是对这三条约束的逐条回应。

分数三种算法：softmax、sigmoid、sqrtsoftplus

三种主流路由方式在 SGLang 里是同一个参数 scoring_func 的三个取值。

softmax：(g) 直接做行内 softmax，取 top-k。softmax 与 top-k 可交换（单调变换），所以既可以先取 top-k 再做 k 个元素的 softmax，也可以先做全行 softmax 再取。Mixtral、Qwen 系列属于这一支。

sigmoid：(\sigma(g_j)) 作为分数，取 top-k，再在选中的 k 个之间归一。DeepSeek-V3 引入的 noaux_tc 路由用这一支。

sqrtsoftplus：(\sqrt{\ln(1+e^{x})})，DeepSeek-V4 系列使用。它的作用是压平尾部：平方根让分数差异收缩，让偏置修正机制在 top-k 边界附近更有效。

三种算法的差异远不只是”激活函数选哪个”。softmax 做全行归一，意味着分数本身已经带上了”所有专家”的信息；而 sigmoid、sqrtsoftplus 是逐元素函数，没有全行归一，所以”归一”必须额外做一次，而且只能在选中的 k 个之间做——这正是 DeepSeek 路由”选完再归一”的结构来源。

入口：TopK 算子与 TopKConfig

SGLang 把路由封装成一个算子类 TopK，继承平台的 BaseFusedOp，按设备分派到 forward_cuda / forward_npu / forward_cpu / forward_xpu / forward_musa。它的构造函数把所有路由相关的配置收进一个 TopKConfig 数据类：

1
# python/sglang/srt/layers/moe/topk.py（节选）
2
@dataclass
3
class TopKConfig:
4
    top_k: int
5
    use_grouped_topk: bool = False
6
    topk_group: Optional[int] = None
7
    num_expert_group: Optional[int] = None
8
    renormalize: bool = True
9
    correction_bias: Optional[torch.Tensor] = None
10
    torch_native: bool = False
11
    routed_scaling_factor: Optional[float] = None
12
    apply_routed_scaling_factor_on_output: bool = False
13
    output_format: Optional[TopKOutputFormat] = None
14
    runner_backend: Optional[MoeRunnerBackendLike] = None
15
    scoring_func: str = "softmax"
16
    sqrtsoftplus_log1p: bool = False
17
    fused_gate_packed_ids: bool = False
18
    allow_routed_experts_capture: bool = True
19
    layer_id: Optional[int] = None

逐项说明：

• top_k：输出的总宽度，包含共享专家的槽位。注释写得很明确：“routed experts plus shared TopK slots”。所以当模型有共享专家且选择融合进 topk 时，top_k 参数本身只是路由专家数，真正的输出宽度是 top_k + num_shared_topk_slots。
• use_grouped_topk / num_expert_group / topk_group：分组路由开关与两个规模参数。
• renormalize：选完后是否在 k 个权重上归一。
• correction_bias：DeepSeek 的 e_score_correction_bias，一个长度为 (N) 的向量；注释里还说明它可以是二维的”逐 token 偏置”（dynamic_bias）。
• apply_routed_scaling_factor_on_output：routed_scaling_factor 是否已经由路由自己乘上。这个区分很重要——后文”缩放折叠”一节讲。
• output_format：路由输出给下游的格式，五种之一，见下文。
• runner_backend：下游 MoE 执行后端（Triton / FlashInfer / AITERT / DeepGEMM 等），它决定 output_format。
• layer_id：全局层号，EPLB 的按层映射、专家分布记录器都按它索引。
• allow_routed_experts_capture：投机解码的草稿层会把它关掉，避免草稿路由污染目标层的专家分布统计。

TopK.__init__ 里还有一处容易被忽略的细节：它从全局的”当前模型”对象读取共享专家配置：

1
shared_experts = get_model().shared_experts
2
self.enable_waterfill = shared_experts.use_waterfill
3
num_shared_slots = shared_experts.num_shared_topk_slots
4
self.topk_config = TopKConfig(
5
    top_k=top_k + num_shared_slots,
6
    ...
7
)

“Waterfill” 是一种按专家容量均衡的路由扩展（标准 topk 之后按剩余容量再分），它要求路由输出是标准格式，因为 Waterfill 需要读 topk_weights / topk_ids 再改写。所以 output_format 在启用 Waterfill 时会强制回到 STANDARD。

选择之前：select_experts 的决策树

真正的分派在 select_experts 里。它的主要分支如下（只保留结构，省略非关键参数）：

1
# python/sglang/srt/layers/moe/topk.py（节选）
2
def select_experts(hidden_states, router_logits, topk_config, *, layer_id=None,
3
                   num_token_non_padded=None, expert_location_dispatch_info=None,
4
                   dynamic_expert_bias=None, router_logits_partials=None):
5
    top_k = topk_config.top_k
6
    use_grouped_topk = topk_config.use_grouped_topk
7
    ...
8
    if dynamic_expert_bias is not None:
9
        topk_weights, topk_ids = biased_grouped_topk_impl(...)      # 逐 token 偏置，永远走 torch 实现
10
    elif use_grouped_topk:
11
        if correction_bias is None:
12
            topk_weights, topk_ids = grouped_topk(...)               # 分组但不带偏置
13
        else:
14
            topk_weights, topk_ids = biased_grouped_topk(...)        # 分组 + 偏置
15
    elif topk_config.torch_native and custom_routing_function is None:
16
        topk_weights, topk_ids = fused_topk_native(...)
17
    elif custom_routing_function is None:
18
        if scoring_func == "sqrtsoftplus" or scoring_func == "sigmoid":
19
            topk_weights, topk_ids = biased_topk_jit_kernel_impl(...)  # JIT 融合路由
20
        elif <flashinfer_trtllm_routed 且 softmax 且无偏置>:
21
            topk_weights, topk_ids = fused_topk_softmax_torch_raw_logits(...)
22
        else:
23
            topk_weights, topk_ids = fused_topk(...)                 # softmax 主路径
24
    else:
25
        topk_weights, topk_ids = custom_routing_function(...)
26
    ...
27
    topk_ids, topk_weights, recorder_topk_ids = _post_process_topk_ids(
28
        topk_ids=topk_ids, topk_weights=topk_weights, topk_config=topk_config,
29
        router_logits=router_logits, layer_id=layer_id,
30
        num_token_non_padded=num_token_non_padded,
31
        expert_location_dispatch_info=expert_location_dispatch_info,
32
        padded_rows_masked=padded_rows_masked,
33
    )
34
    return StandardTopKOutput(topk_weights, topk_ids, router_logits)

决策树的要点：

• 按模型配置分支，不按运行时形状分支。路由形式由模型结构决定，启动之后就固定。唯一随形状变化的是 torch_native（纯 torch 实现），它只在硬件没有专用 kernel、或 torch.compile 有收益（batch size 1）时开启。
• dynamic_expert_bias 总是退回 torch 实现。逐 token 偏置意味着偏置表形状是 [T, N]，现有 JIT kernel 只接受”一组向量”；这一点后文单独讲。
• 分组 + 偏置走一条独立分支：注意 grouped_topk（无偏置）与 biased_grouped_topk（有偏置）是两个函数，后者内部再选 JIT / FlashInfer / torch 实现。

分组路由的实现：先把组选出来，再在组里选专家

DeepSeek-V3 的每个 MoE 层有 1 个共享专家、256 个路由专家，路由专家分成 8 组、每组 32 个；每 token 选 8 个路由专家，共享专家恒被激活。SGLang 的路由层要处理的就是这张图里”256 选 8、8 组选 4 组”的部分，再把共享专家以额外槽位的方式拼进输出。

grouped_topk_gpu 和 biased_grouped_topk_impl 是理解 DeepSeek 路由的两份”可读参考实现”。前者不带偏置，用 torch.compile 编译；后者带偏置，是 JIT kernel 不满足条件时的兜底。

先看无偏置版本的核心：

1
# python/sglang/srt/layers/moe/topk.py（节选，grouped_topk_gpu）
2
scores = torch.softmax(gating_output, dim=-1) if scoring_func == "softmax" else gating_output.sigmoid()
3
num_token, num_experts = scores.shape
4
group_scores = scores.view(num_token, num_expert_group, -1).max(dim=-1).values   # [T, G]
5
group_idx = torch.topk(group_scores, k=topk_group, dim=-1, sorted=False)[1]       # [T, k_g]
6
group_mask = torch.zeros_like(group_scores)
7
group_mask.scatter_(1, group_idx, 1)
8
score_mask = group_mask.unsqueeze(-1).expand(num_token, num_expert_group,
9
              scores.shape[-1] // num_expert_group).reshape(num_token, -1)      # [T, N]
10
tmp_scores = scores.masked_fill(~score_mask.bool(), 0.0)
11
topk_weights, topk_ids = torch.topk(tmp_scores, k=topk, dim=-1)

这段的语义是：每组取组内最大分作为组代表分，选前 topk_group 组，把其余组的专家分数清零，再在全行取 top-k。注意”清零”而不是设成负无穷：masked_fill(..., 0.0) 对 sigmoid/softmax 分数是安全的，因为分数非负，清零后不可能进入 top-k。但若本批 slot 总数少于专家数（例如 batch size 1、k = 2 时只有 2 个 slot），那么”可选的专家”只有 2 个，剩下 k − 2 个槽位必然落到被清零的 0 分专家上——下游就会读到”得分 0 的专家 id”。因此这条路径主要作为参考实现，以及 FlashInfer fused_topk_deepseek 与 JIT 路由都不可用的兜底；生产路径走的是下文的下标路线，padding 行有单独的 -1 约定。

再看真正在 DeepSeek-V3 上用的带偏置版本：

1
# python/sglang/srt/layers/moe/topk.py（节选，biased_grouped_topk_impl）
2
scores = gating_output.sigmoid()
3
scores_for_choice = scores.view(num_token, -1) + correction_bias.view(-1, num_experts)
4
group_scores = (scores_for_choice.view(num_token, num_expert_group, -1)
5
                .topk(2, dim=-1)[0].sum(dim=-1))                                  # [T, G]
6
group_idx = torch.topk(group_scores, k=topk_group, dim=-1, sorted=False)[1]
7
group_mask = torch.zeros_like(group_scores)
8
group_mask.scatter_(1, group_idx, 1)
9
score_mask = ...
10
tmp_scores = scores_for_choice.masked_fill(~score_mask.bool(), float("-inf"))
11
_, topk_ids = torch.topk(tmp_scores, k=topk, dim=-1)
12
topk_weights = scores.gather(1, topk_ids)      # 权重取自未加偏置的 scores

与无偏置版本的三个差别：

1. 组代表分是组内 top-2 的和，而不是最大值。DeepSeek-V3 的技术报告里这个设计让一个”单点极强但整体偏弱”的组更难挤进前 topk_group，让组的选择更看整体；
2. 未选中的组用 -inf 屏蔽而不是 0，与后面选专家时”任何未被组选中的专家不可能被选中”的语义一致；
3. topk_weights = scores.gather(1, topk_ids)：用未加偏置的 scores 取权重。这是全篇最重要的语义点——偏置只改”选谁”，不改”算多少”。如果偏置也进入权重，负载均衡的力度就会直接扰动模型输出；训练时 e_score_correction_bias 作为独立的偏置参数（只作用于选择阶段）才能实现”不依赖辅助损失的负载均衡”（noaux_tc）。

这套机制来自字节跳动 Seed 团队的 Loss-Free Balancing 工作（论文，2024 年 8 月）。它的要点正如图 2 所示：每个训练 step 按”带偏置的门控分数”选专家，step 结束后再按各专家的负载更新偏置——负载超载的专家偏置下调，闲置的专家偏置上调。

推理阶段偏置已经固定，本篇文章关心的是它被”怎样使用”：只进入选择分数的比较，不进入选中专家的权重。上图画的是训练时偏置的更新循环，推理时只剩”用偏置选、用原分算”这半圈。

同一个文件里 grouped_topk_gpu 选组用的是 max，而 biased_grouped_topk_impl 用组内 top-2 之和，这是两个不同路由配置的差别，不是实现不一致——前者的配置来自 topk_method == "greedy" 那类模型，后者来自 "noaux_tc"。

生产路径：JIT Triton 路由内核逐段拆

真正在生产中跑的是 python/sglang/kernels/ops/moe/moe_fused_gate.py 里的 _router_triton_kernel。它用一行 token 分块（BLOCK_M 行）乘一个专家分块（BLOCK_N，N 向上的 2 的幂），在寄存器与共享内存里把”打分 → 修正 → 选组 → 选专家 → 归一 → 写出”全部完成。下面按段看。

1
# python/sglang/kernels/ops/moe/moe_fused_gate.py（节选）
2
pid = tl.program_id(0)
3
offs_m = pid * BLOCK_M + tl.arange(0, BLOCK_M)
4
offs_n = tl.arange(0, BLOCK_N)
5
mask_m = offs_m < M
6
mask_n = offs_n < N
7
if HAS_BIAS:
8
    bias = tl.load(bias_ptr + offs_n * stride_bias, mask=mask_n, other=0.0).to(tl.float32)
9
else:
10
    bias = tl.zeros([BLOCK_N], dtype=tl.float32)
11
...
12
scores = tl.load(row_ptr, mask=mask2d, other=0.0).to(tl.float32)

1
if SCORING_FUNC == 0:
2
    activated = tl.sigmoid(scores)
3
    biased = activated + row_bias
4
elif SCORING_FUNC == 1:
5
    sp = tl.where(scores > 20.0, scores, libdevice.log1p(libdevice.exp(scores)))
6
    activated = libdevice.sqrt(sp)
7
    biased = activated + row_bias
8
else:
9
    logit = scores
10
    if HAS_SOFTCAP:
11
        z = logit / moe_softcapping
12
        logit = moe_softcapping * (2.0 * tl.sigmoid(2.0 * z) - 1.0)   # tanh 的 sigmoid 形式
13
    biased = logit + row_bias
14
    row_max = tl.max(biased, axis=1)[:, None]
15
    exp_row = tl.where(mask_n[None, :], tl.exp(biased - row_max), 0.0)
16
    activated = exp_row / tl.sum(exp_row, axis=1)[:, None]

1
if N_GROUP > 1:
2
    offs_g = tl.arange(0, BLOCK_G)
3
    group_of_n = offs_n // EXPERTS_PER_GROUP
4
    group_score = tl.full([BLOCK_M, BLOCK_G], -float("inf"), dtype=tl.float32)
5
    for g in tl.static_range(N_GROUP):
6
        in_g = (group_of_n[None, :] == g) & mask_n[None, :]
7
        vals = tl.where(in_g, biased, -float("inf"))
8
        top1 = tl.max(vals, axis=1)[:, None]
9
        vals2 = tl.where(vals >= top1, -float("inf"), vals)
10
        top2 = tl.max(vals2, axis=1)[:, None]
11
        group_score = tl.where(offs_g[None, :] == g, top1 + top2, group_score)

1
    gcur = group_score
2
    keep = tl.zeros([BLOCK_M, BLOCK_N], dtype=tl.float32)
3
    for _i in tl.static_range(TOPK_GROUP):
4
        gmax = tl.max(gcur, axis=1)[:, None]
5
        glane = tl.where(gcur == gmax, offs_g[None, :], N_GROUP + 1)
6
        win_g = tl.min(glane, axis=1)[:, None]          # 平局取组号最小的
7
        keep = tl.where(group_of_n[None, :] == win_g, 1.0, keep)
8
        gcur = tl.where(offs_g[None, :] == win_g, -float("inf"), gcur)
9
    biased = tl.where(keep > 0.0, biased, -float("inf"))

1
    cur = biased
2
    remaining = tl.broadcast_to(mask_n[None, :], (BLOCK_M, BLOCK_N))
3
    for k in tl.static_range(K_ROUTED):
4
        max_val = tl.max(cur, axis=1)[:, None]
5
        is_max = remaining & (cur == max_val)
6
        lane_id = tl.where(is_max, offs_n[None, :], N + 1)
7
        win_lane = tl.min(lane_id, axis=1)[:, None].to(tl.int32)
8
        win_activated = tl.sum(tl.where(offs_n[None, :] == win_lane, activated, 0.0), axis=1)[:, None]
9
        slot = offs_k[None, :] == k
10
        selected_vals = tl.where(slot, win_activated, selected_vals)
11
        selected_idx = tl.where(slot, win_lane, selected_idx)
12
        remaining = remaining & (offs_n[None, :] != win_lane)
13
        cur = tl.where(remaining, cur, -float("inf"))

1
    routed_sum = tl.sum(tl.where(mask_k_routed[None, :], selected_vals, 0.0), axis=1)[:, None]
2
    if RENORMALIZE:
3
        if RENORMALIZE_EPSILON > 0.0:
4
            norm = routed_sum + RENORMALIZE_EPSILON
5
        else:
6
            norm = tl.where(routed_sum > 0.0, routed_sum, 1.0)
7
        selected_vals = selected_vals / norm
8
    if APPLY_SCALE:
9
        selected_vals = selected_vals * routed_scaling_factor
10
    if K_ROUTED < K:
11
        is_shared = (offs_k[None, :] >= K_ROUTED) & mask_k_total[None, :]
12
        shared_weight = 1.0 if APPLY_SCALE else 1.0 / routed_scaling_factor
13
        shared_idx = (N + (offs_k - K_ROUTED)).to(tl.int32)[None, :]
14
        selected_vals = tl.where(is_shared, shared_weight, selected_vals)
15
        selected_idx = tl.where(is_shared, shared_idx, selected_idx)

1
    if HAS_PADDING:
2
        selected_vals = tl.where(live_m[:, None], selected_vals, 0.0)
3
        selected_idx = tl.where(live_m[:, None], selected_idx, -1)

1
    if HAS_PACKED:
2
        w_bits = selected_vals.to(tl.bfloat16).to(tl.int16, bitcast=True).to(tl.int32)
3
        packed = (selected_idx << 16) | (w_bits & 0xFFFF)

宽行专家：从 k 轮 argmax 到 Radix-Select

主线内核的 K_ROUTED 轮顺序 argmax 有一个隐含成本：每一轮都要把长度为 BLOCK_N 的向量做一次全行归约。Mixtral 的 8 个专家、DeepSeek-V3 的 256 个专家都没问题；但 Kimi K3 这类模型有 896 个专家、top-16，16 轮 1024 宽度的归约开始主导路由耗时。

SGLang 的对策是在 moe_fused_gate 入口加一条专用分支：满足 scores 宽度 896、topk == 16、bias 为 fp32、行连续、dtype 为 bf16/fp32 时，调用 moe_route_radix.route_radix。这个 Radix-Select 内核的工作方式与 argmax 循环完全不同：

• 键留在寄存器里：224 个线程、每个线程负责 4 个专家，把 896 个分数放进寄存器，不做全局排序；
• 按字节做基数选择：把 fp32/bf16 的分数按二进制位分成”桶”，每一轮只看一个字节（8 位），从高位到低位逐字节收窄候选集合；某一轮若 top-k 已经能在该字节边界上完整分出，后续低位直接跳过（注释里的 “rounds exit early when the top-k separates on a byte boundary”）；
• 前缀和用 warp scan 而不是 CUB：每个 warp 内部的桶计数用 shuffle 前缀和，避免为小规模数据启动 CUB 的 block scan；
• 输出顺序可省：主线的 k 轮 argmax 返回的是”按分数降序、平局按 id 升序”；Radix 内核默认 sorted=False，返回”按专家 id 升序”的结果。下游 GEMM / combine 对槽位顺序不敏感，因此跳过最后的排序。

代码注释给出的实测结论是：在 B200 上，[1..8192, 896] top-16 区间内比 Triton 主线快 3.1–3.5 倍；两条路径的获胜集合逐位一致，只有归一化求和的顺序不同，权重最多差约 1 ulp。这个”集合一致、顺序不同”的取舍值得记下来：它说明路由层的正确性标准是”选出的专家集合 + 每个专家的权重”（权重对求和顺序不敏感），而不是”输出槽位的排列顺序”。

这也解释了为什么 Radix 路径要求那么严格：num_fused_shared_experts == 0、num_expert_group <= 1、无 softcap、无 padding、无 packed 输出。任何一个条件不满足，内核就必须维持”排序后的槽位顺序”或额外处理共享列，省排序的收益就消失了。SGLang 选择”只在收益确定时启用”，而不是为了覆盖率把 Radix 做成通用路径。

共享专家：拼进 topk，还是单独走一条路

共享专家在 SGLang 里不是”特例分支”，而是一组被追加在路由专家之后的固定槽位。TopKConfig.num_shared_topk_slots 决定追加几个，_post_process_topk_ids 负责把它们拼到 topk_ids / topk_weights 的最后一列。拼接方式有两种：

1. 全局共享槽位：所有 rank 共享同一组共享专家 id，权重恒为 1（或 (1/s)，见缩放一节）。适合共享专家被复制到每个 rank 的部署。
2. 按 rank 的共享槽位（has_per_rank_fused_shared_slots）：每个 rank 有自己的共享专家副本，remap_topk_for_per_rank_shared_slots 会把路由 id 先按”每 rank 的本地路由专家数”重排，再在对应 rank 的块后插入共享槽位。这个函数里最容易被忽略的是：

1
# python/sglang/srt/layers/moe/topk.py（节选）
2
routed_ids = topk_ids + (topk_ids // num_local_routed) * num_fused_shared_experts
3
shared_ids = ep_rank * num_local_experts + num_local_routed + torch.arange(num_fused_shared_experts, ...)

(topk_ids // num_local_routed) * S 这一项是”每经过一个 rank 的专家块，就为它前面的共享槽位留出 S 个位置”。如果只加一个固定偏移（+ S），当 S > 1 时后一个 rank 的路由 id 会与前一个 rank 的共享槽位撞号。这个细节说明：共享专家一旦进入 topk_ids，它就成为 EP 分块布局的一部分，必须参与逻辑到物理的地址计算，不能再当成”额外的一条稠密路径”。

与 DeepInfer 的对照很能说明两种设计的取舍。DeepInfer 把共享专家实现成 ReplicatedLinear，在 MoE 分支之外单独算一条稠密 MLP；SGLang 把共享专家拼成额外 topk 槽位。前者省掉了”每个 token 多一个路由槽位”的分桶与 GEMM 开销，但多一次独立的小矩阵乘与一次加法；后者让共享专家与路由专家共用同一套 grouped GEMM，代价是分桶数组多一列、EP 布局需要重排。哪种更好取决于 decode batch 大小：小 batch 下独立 MLP 的 launch 开销更贵，大 batch 下共享槽位的分桶浪费更明显。

五种输出格式：为什么下游要看到不同的账本

TopKOutputFormat 定义了五种格式，对应不同后端对”路由结果”的不同消费方式：

这里最值得讲的是 BYPASSED。FlashInfer 的 TRT-LLM 风格路由 MoE 内核（从 TRT-LLM 移植过来）会在自己的 kernel 里做打分、top-k 与分桶，如果 SGLang 再前置一个 select_experts，就是”路由两次”。所以 TopK 在 bind_runner_backend 时检测到后端是 flashinfer_trtllm 且路由配置满足 _can_bypass_trtllm_routing，就把输出格式设为 BYPASSED，并把原始 logits 原样交出去：

1
# python/sglang/srt/layers/moe/topk.py（节选，TopK.forward_cuda）
2
elif output_format == TopKOutputFormat.BYPASSED:
3
    return BypassedTopKOutput(
4
        hidden_states=hidden_states,
5
        raw_router_logits=router_logits,
6
        topk_config=self.topk_config,
7
        num_token_non_padded=num_token_non_padded,
8
        expert_location_dispatch_info=expert_location_dispatch_info if requires_remapping else None,
9
    )

BypassedTopKOutput.to_standard() 则提供了”按需物化”的出口：如果下游某个 kernel 不支持内部路由，就调用 to_standard() 再走一遍 select_experts。而它里面的 expert_location_dispatch_info 一旦非空（EPLB 需要逻辑到物理重映射），就会在 router_logits 属性访问时抛 NotImplementedError——注释说得很直接：“Expert placement remapping requires materialized TopK routing; this MoE kernel routes directly from logits.” 也就是说，一旦启用专家重映射（动态 EPLB），就必须把路由物化成 topk_ids，因为下游内核自己算出来的逻辑专家号没有办法被”路由之外”的层改写。

这条约束把本篇后半程的逻辑串起来了：EPLB 越积极，路由输出越必须显式。

后处理：_post_process_topk_ids 里的四件事

不管走哪条内核，select_experts 最后都会把 topk_ids / topk_weights 交给 _post_process_topk_ids 处理。从源码看，它做的是四件事：共享列追加、逻辑到物理重映射、padding 行掩码、专家分布记录。

1
# python/sglang/srt/eplb/expert_location_dispatch.py（节选）
2
def topk_ids_logical_to_physical(topk_ids, info, log2phy_prob=None):
3
    if info is None:
4
        return topk_ids
5
    if info.ep_dispatch_algorithm == "static":
6
        return _topk_ids_logical_to_physical_static(topk_ids, info)
7
    if info.ep_dispatch_algorithm in ["dynamic", "fake"]:
8
        return _topk_ids_logical_to_physical_dynamic(topk_ids, info)
9
    if info.ep_dispatch_algorithm == "lp":
10
        return _topk_ids_logical_to_physical_probability(topk_ids, info, log2phy_prob)
11
    raise NotImplementedError(...)
12

13
def _topk_ids_logical_to_physical_static(topk_ids, info):
14
    physical_topk_ids = info.partial_logical_to_rank_dispatch_physical_map[topk_ids]
15
    return physical_topk_ids.to(topk_ids.dtype)

1
# python/sglang/srt/layers/moe/topk.py（节选）
2
def _mask_topk_ids_padded_region(topk_ids, num_token_non_padded=None, fill_value=-1):
3
    if num_token_non_padded is None:
4
        return
5
    ...
6
    if _is_cuda and topk_ids.dtype == torch.int32 and fill_value == -1:
7
        mask_topk_ids(topk_ids, num_token_non_padded)

1
    if recorder_topk_ids is not None:
2
        get_global_expert_distribution_recorder().on_select_experts(topk_ids=recorder_topk_ids)

从 topk 到分组 GEMM：moe_align_block_size 的排序账

topk_ids 是 [T, k] 的二维表，分组 GEMM 要的是”按专家连续排布、每个专家一段、每段长度对齐到 block_size”。中间的转换函数是 moe_align_block_size，定义在 python/sglang/srt/layers/moe/moe_runner/triton_utils/moe_align_block_size.py。它的 docstring 给出一个可以直接手算的例子：

Given topk_ids = [[2, 3, 4], [1, 2, 4], [1, 3, 4], [1, 2, 3]], block_size = 4, num_experts = 4:

• 12 (token, expert) slots, 4 experts, 3 tokens each;
• pad 1 token per expert since block_size is 4;
• flatten → [2,3,4, 1,2,4, 1,3,4, 1,2,3]; append padding [12,12,12,12] per block;
• sorted by expert → [3,6,9,12, 0,4,10,12, 1,7,11,12, 2,5,8,12]; tokens 12 are padding.

这 16 个数字就是 sorted_token_ids；配一个 expert_ids = [0,1,2,3]，num_tokens_post_padded = 16，一次 4 块 × 每块 4 行的 GEMM 就正好对应”4 个专家 × 每专家 4 个 slot”。

图中把 topk_ids 的每个 (token, slot) 压成一个全局槽位号，再按 expert + 1 分桶。真实槽位用蓝色，补齐用的空洞用灰色：空洞里写的不是 0，而是 numel = T × k，即”所有真实 token 之后的下标”。下游 GEMM 读到 numel 就知道这一行是填充，按 0 权重丢弃；这样 padding 永远不会与真实 token 0 撞号。

这里有一个初学者最容易写错的细节：padding 值为什么是 numel 而不是 0。如果 padding 写 0，那么一个”专家 A 的第 1 个 slot”和”专家 A 的补齐行”都会指向 token 0，GEMM 会把 token 0 的隐状态算两遍，再把两份结果都写回 token 0，导致输出翻倍。写 numel 之后，填充行的 token 下标越界，GEMM 要么用 mask 跳过，要么读到一个不会被写回的位置；总之不会污染真实 token。这个约定在 vLLM 的 moe_align_block_size 里同样存在，是 MoE 推理内核的公共常识。

先看容量预算：

1
# python/sglang/srt/layers/moe/moe_runner/triton_utils/moe_align_block_size.py（节选）
2
if topk_ids.numel() < num_experts + 1:
3
    max_num_tokens_padded = topk_ids.numel() * block_size
4
else:
5
    max_num_tokens_padded = topk_ids.numel() + (num_experts + 1) * (block_size - 1)
6
max_num_tokens_padded = (max_num_tokens_padded + 3) // 4 * 4       # int4 写入对齐
7
sorted_ids = torch.empty((max_num_tokens_padded,), dtype=torch.int32, device=topk_ids.device)
8
max_num_m_blocks = triton.cdiv(max_num_tokens_padded, block_size)
9
expert_ids = torch.empty((max_num_m_blocks,), dtype=torch.int32, device=topk_ids.device)
10
num_tokens_post_pad = torch.empty((1,), dtype=torch.int32, device=topk_ids.device)
11
cumsum_buffer = torch.empty((num_experts + 2,), dtype=torch.int32, device=topk_ids.device)

max_num_tokens_padded 的推导值得一算：

• 如果本批 slot 总数少（numel < E + 1），每个专家最多装 numel 个，对齐后是 numel × block_size，这是最坏情况（所有 slot 挤一个专家）；
• 否则，假设 (E + 1) 个桶（含无效桶）都至少有一个 slot，每个桶的最后一块最多补 block_size - 1 个，所以上界是 numel + (E + 1) × (block_size - 1)。

后者精确到只按”真正有 token 的桶”计算会更紧，但 kernel 在 GPU 上无法方便地先数”有几个桶非空”再分配，所以用这个上界——多分配的一点显存远小于”分配不足”带来的隐患。(x + 3) // 4 * 4 是为了让后面用 int4 向量写入 padding。

cumsum_buffer 长度为 E + 2：topk_ids + 1 之后的桶号范围是 [0, E]（-1 → 0，E-1 → E），再加一个位置给前缀和。

1
// python/sglang/kernels/jit/csrc/moe/moe_align_kernel.cu（节选）
2
if (blockIdx.x == 1) {
3
  if (pad_sorted_token_ids) {
4
    Vec fill_vec;
5
    fill_vec.x = fill_vec.y = fill_vec.z = fill_vec.w = numel;
6
    int32_t total_vecs = max_num_tokens_padded / VEC_SIZE;
7
    Vec* out_ptr = reinterpret_cast<Vec*>(sorted_token_ids);
8
    for (int32_t i = threadIdx.x; i < total_vecs; i += blockDim.x) out_ptr[i] = fill_vec;
9
    for (int32_t i = total_vecs * VEC_SIZE + threadIdx.x; i < max_num_tokens_padded; i += blockDim.x)
10
      sorted_token_ids[i] = numel;
11
  }
12
  return;
13
}

1
  extern __shared__ int32_t smem[];
2
  int32_t* shared_counts = smem;
3
  int32_t* prefix = shared_counts + num_experts;
4
  int32_t* scan_buf = prefix + num_experts + 1;
5
  __shared__ int32_t s_total_tokens_post_pad;
6
  ...
7
  for (size_t i = tid; i < numel; i += stride) {
8
    int expert_id;
9
    if constexpr (kMapInvalidExpertsToSentinel) expert_id = sentinel_expert_bucket(topk_ids[i], num_experts);
10
    else expert_id = topk_ids[i] + 1;
11
    atomicAdd(&shared_counts[expert_id], 1);
12
  }
13
  __syncthreads();
14
  int32_t padded_count = 0;
15
  if (tid < num_experts) {
16
    int32_t count = shared_counts[tid];
17
    padded_count = (count + block_size - 1) / block_size * block_size;
18
    scan_buf[tid] = padded_count;
19
  }

1
template <typename scalar_t>
2
__device__ __forceinline__ int32_t sentinel_expert_bucket(scalar_t expert_id, int32_t num_buckets) {
3
  using unsigned_t = std::make_unsigned_t<scalar_t>;
4
  const unsigned_t shifted = static_cast<unsigned_t>(expert_id) + unsigned_t{1};
5
  return shifted < static_cast<unsigned_t>(num_buckets) ? static_cast<int32_t>(shifted) : 0;
6
}

1
  if (tid <= num_experts) cumsum[tid] = prefix[tid];
2
  const int32_t num_blocks = s_total_tokens_post_pad / block_size;
3
  for (int32_t i = tid; i < num_blocks; i += stride) {
4
    int32_t block_start = i * block_size;
5
    int left = 0, right = num_experts;
6
    while (left < right) {
7
      int mid = (left + right) >> 1;
8
      if (prefix[mid] <= block_start) left = mid + 1;
9
      else right = mid;
10
    }
11
    expert_ids[i] = left - 2;
12
  }
13
  if constexpr (kPadExpertIds) {
14
    for (int32_t i = num_blocks + tid; i < max_num_m_blocks; i += stride) expert_ids[i] = -1;
15
  }

1
if (_is_cuda and topk_ids.numel() <= SMALL_NUMEL_LIMIT
2
        and num_experts + 1 > _CUDA_SMALL_BATCH_MAX_BUCKETS
3
        and not ignore_invalid_expert and not pad_expert_ids):
4
    moe_align_small_numel(...)
5
    return sorted_ids, expert_ids, num_tokens_post_pad

缩放系数放在哪里：一次折叠省一个 kernel

routed_scaling_factor（DeepSeek-V3 是 2.5）在数学上只是一个乘法，但它和”共享专家权重恒为 1”、“归一是否在 k 个上进行”、“下游 epilogue 会不会再乘一次”纠缠在一起。SGLang 用两件事把它理清楚：

1. apply_routed_scaling_factor_on_output：告诉下游”路由已经把缩放乘进去了，epilogue 不要再乘”；
2. fold_scale_after_routing：某些路径（torch native、自定义路由、非 JIT 内核）的内核不接受缩放参数，此时 select_experts 在拿到输出后自己补乘：

1
# python/sglang/srt/layers/moe/topk.py（节选）
2
if (fold_scale_after_routing and apply_routed_scaling_factor_on_output
3
        and routed_scaling_factor is not None):
4
    num_routed_for_gate = top_k_for_gate - num_shared_topk_slots_for_gate
5
    topk_weights[:, :num_routed_for_gate] *= routed_scaling_factor

注意只乘 [:num_routed_for_gate]：共享列保持自己的权重（1 或者 (1/s)），不会被这次补乘再改一遍。

为什么”内核里乘”比”输出后再乘”好？decode 时 topk_weights 是 [T, k]、T 可能只有几十，这个乘法本身可以忽略；真正要省的是一次额外的 kernel launch。在每层都要做一次乘法时，几十层的 decode 会产生几十次 launch，每次 launch 即使只有几微秒，也积少成多。把缩放写进路由内核的 APPLY_SCALE 编译期开关，就等于在已有的 tl.store 之前多一条 mul 指令，零额外 launch。

同理，is_routed_scaling_factor_fused_in_topk 这个函数存在的意义是”让下游知道要不要跳过缩放”：

1
def is_routed_scaling_factor_fused_in_topk(topk_config):
2
    if topk_config.apply_routed_scaling_factor_on_output:
3
        return True
4
    if not _use_aiter:
5
        return False
6
    if topk_config.use_grouped_topk:
7
        return topk_config.correction_bias is not None
8
    return (topk_config.scoring_func == "sqrtsoftplus"
9
            and topk_config.custom_routing_function is None)

这段说明 AITER 后端”总是在内核里折叠缩放”，而其它后端要显式传参。跨后端一致性靠这一个函数显式维护，而不是靠”大家都记着别乘两次”。

逐 token 偏置与”fake 路由”：工程细节两则

1
if dynamic_expert_bias is not None:
2
    if scoring_func != "sigmoid":
3
        raise ValueError("Per-token expert bias is only supported with sigmoid routing")
4
    topk_weights, topk_ids = biased_grouped_topk_impl(...)

1
def transform_select_experts_inputs(router_logits, correction_bias, info):
2
    if (info is not None) and (info.ep_dispatch_algorithm == "fake"):
3
        router_logits.uniform_(5, 10)
4
        if correction_bias is not None:
5
            correction_bias = torch.zeros_like(correction_bias)
6
    return router_logits, correction_bias

精度账：为什么 logit 是 fp32、ids 是 int32

路由层是 MoE 前向里精度要求最微妙的一段。几个看起来”无所谓”的决定都有原因：

• logit 强制 fp32。TopK.__init__ 说明门控层的输出是 fp32（例如 ReplicatedLinear(..., use_fp32_output=True)），JIT kernel 的注释也写明 “raw logits, fp32/fp16/bf16 (upcast to fp32 on load)“。原因是 top-k 本质是排序，FP16 只给 10 位尾数，两个接近的专家分可能被抹平成相等，使平局分辨率依赖归约顺序；fp32 让”非常接近但不相等的两个分数”仍能区分。
• 累加在 fp32。routed_sum 的求和与除法都在 fp32 下进行；torch 参考实现里 topk_weights.sum(dim=-1, dtype=torch.float32) 也显式指定了 fp32 累加。
• 归一用”零和保护”而不是 epsilon。JIT kernel 的默认 RENORMALIZE_EPSILON = 0.0，走 norm = tl.where(routed_sum > 0.0, routed_sum, 1.0)；torch 路径用 + 1e-20。两者对”和为 0 时”的语义一致（把分母换成 1，权重保持为原值），差异只在极小正和时的精度——1e-20 在 fp32 的可表示范围内（最小正规数约 (1.18 \times 10^{-38})），不会下溢为 0。
• topk_ids 为 int32。一是 moe_align 的桶号需要 E + 1，而 E = 896 尚能塞进 int16；但 padding 用 -1 以及 sentinel 映射用到无符号解释，窄类型反而增加边界情况。二是打包格式 (id << 16) | weight_bits 本身要求 int32。三是下游 Triton GEMM 的索引运算习惯 int32，避免每次读 id 都做一次窄转宽。

精度的一般背景可参考 答疑特辑（二）：浮点数五问。

分组路由的出处：DeepSeekMoE 与 Loss-Free Balancing

本节把”分组选组 + 偏置修正”这套机制的来源补上，避免把 SGLang 的实现当成凭空设计。

分组本身就是 DeepSeekMoE 的结构（DeepSeekMoE 完全拆解）。把 256 个专家切成 8 组、每组 32 个，是”细粒度专家切分”的自然结果：专家越细，单次 top-k 选出的 8 个专家越可能扎堆在同一组里，路由容量被浪费在”一个组的不同专家”身上。分组选组强制先保证每组至少有一个代表，相当于在”专家多样性”上加一道软约束。图 1 里右侧”细粒度切分 + 共享专家”那一列，就是分组路由的硬件前提。

从左到右：传统 MoE 的专家很大、知识混杂；细粒度切分把专家拆小、让路由更精确；共享专家再把”公共知识”从路由路径里摘出来，保证所有 token 都能看到它。SGLang 的 num_fused_shared_experts 和分组 topk，本质上是在执行这张图里的后两步。

修正偏置来自 Loss-Free Balancing（论文，字节跳动 Seed 团队，2024 年 8 月）。传统方法用辅助损失强迫负载均衡，辅助损失越大越均衡、但越伤模型质量；Loss-Free Balancing 改用一组”每个专家一个”的偏置向量 (b_j)：训练时每个 step 结束后，负载超高的专家下调 (b_j)、负载过低的专家上调，路由只拿 (s_j + b_j) 去比较。由于 (b_j) 不参与权重计算，它对模型输出的影响被限制在”改变谁来处理这个 token”这一层，而不是”改变每个专家贡献多少”。SGLang 的 correction_bias 参数正是把这个训练期学到的向量带进推理。论文的 Load-Free Balancing 图示（Figure 1）画出了这张偏置更新循环，SGLang 在推理时只执行它的右半圈——“用偏置选”。

为什么 SGLang 还要保留 torch 参考实现：grouped_topk_gpu 与 biased_grouped_topk_impl 都是 torch.compile 或纯 torch 版本。它们不是”性能路径”，而是”语义基线”。当 JIT kernel 在某些特殊配置（逐 token 偏置、专家数不是 2 的幂、topk_group 与 num_expert_group 的组合不被 FlashInfer 接受）下降级到 torch 实现时，两者必须给出同一组 topk_ids。这一点决定了 SGLang 在写 JIT kernel 时要严格遵守 torch.topk 的平局语义（下标较小者优先），否则两条路径会在同一 logit 上选出不同专家。

排查清单：路由层最容易出错的八个地方

把这一层读完之后，遇到”MoE 输出不对 / 吞吐异常 / 结果随 batch 漂移”时，可以按下面顺序排查。每一条都对应源码里一个具体的约定。

一、逻辑专家号与物理槽位混用。 topk.py 产出的是物理 id，select_experts 入口收到的 router_logits 宽度是逻辑专家数。二者在 ep_dispatch_algorithm is None 时相等，一旦启用 EPLB 就不再相等。把 topk_ids 直接当逻辑专家用于统计，会把”物理槽位 7”误认为”逻辑专家 7”。

二、padding 行的 -1 被当成真实专家。 _mask_topk_ids_padded_region 在 CUDA 上写 -1，HIP 上写 0；两种约定都要求下游 moe_align_block_size 能识别。若自己实现下游内核而没有使用 ignore_invalid_expert，-1 + 1 = 0 会被当成”0 号专家”，导致 padding 行参与真实 GEMM。

三、共享专家的缩放被乘了两次。 若路由层已经把 routed_scaling_factor 乘进 topk_weights，下游 epilogue 不能再乘；反之，若共享权重被预写成 (1/s)，而下游又跳过了缩放，共享专家就只剩 (1/s) 的贡献。判断依据是 is_routed_scaling_factor_fused_in_topk，不要凭”我记得这里乘过”。

四、归一化的范围包含了共享列。 路由专家的权重归一只能在路由专家之间做；把共享列算进分母会让共享专家的固定权重被压低（或抬高）。主线内核用 routed_sum 显式排除共享列，torch 参考实现用 topk_weights[:, :-num_shared] 求和，两处是同一个语义。

五、分组选组用最大值而非 top-2 之和。 grouped_topk_gpu（无偏置）用 max，biased_grouped_topk_impl（有偏置）用组内 top-2 之和。把两者的组分数混用，会让”组选择”在最边界的情况下出现差异。选择哪种取决于模型的 topk_method，不是性能开关。

六、平局规则不一致。 JIT 内核在”选组”和”选专家”两处都用”id 最小者胜”；torch 的 torch.topk 在 sorted=False 时也保证下标较小者优先。自定义路由函数（custom_routing_function）如果用了不稳定的排序，会让同一请求在不同后端产生不同的 topk_ids。排查”同一 prompt 在 A/B 两个后端结果不同”时，先查这里。

七、output_format 与下游后端不匹配。 BYPASSED 格式意味着路由没有物化；一旦启用 EPLB 重映射，BypassedTopKOutput.router_logits 会直接抛异常。这不是 bug，而是设计约束。看到这个异常时，要么关掉动态 EPLB，要么把后端换成消费 STANDARD 的 kernel。

八、KV 页大小与 MoE 块大小混淆。 moe_align_block_size 的 block_size 是”分组 GEMM 的行块”，与 PagedAttention 的 KV page size（SGLang 默认 1 或 128 token）完全无关。把两者当成同一个参数，会在调优时改错地方：前者影响 GEMM 的对齐浪费，后者影响 KV 的碎片率。

与 DeepInfer 和 vLLM 的横向对照

同一条”token → 专家”的路，在三套框架里被安排在完全不同的位置。

DeepInfer（DeepInfer 源码拆解（三十五）：MoE 层的 MLP）把门控、选专家、逻辑到物理映射一起塞进了同一个内核（moe_topk_gate，或 MegaMoE 路径下的 bf16_mega_gate）。它这样做的原因是：标准路径上门控输出 [T, 384] 的 fp32 logits 要先落盘再读回，在 decode 的小 T 下这次中间读写与 GEMM 本身相当。融合之后，logits 只存在于寄存器/共享内存。代价是”路由逻辑”与”专家执行逻辑”绑定得比较紧，换后端就要动内核。

vLLM 的 fused_topk 是 Triton 或 CUDA 内核，但把分组路由、偏置、映射拆成多个较小的步骤，并允许不同 MoE 后端（Triton、CUTLASS、Marlin、DeepGEMM）各自消费 topk_ids。它对”逻辑到物理”的处理依赖 expert_map，但分桶一般在后端内部做。

SGLang 处于中间位置：路由内核本身是融合的（一次 launch 产出权重与 id），但”逻辑到物理”、“共享列追加”、“padding 掩码”、“分布记录”留在 Python 层，并通过 TopKOutputFormat 决定”要不要真的物化”。它的灵活性主要来自 BYPASSED——对于”有能力自己路由”的后端，SGLang 不重复劳动，只做 EPLB 校验。

三者面对同一组硬约束（不能同步主机、账本要能直接下标、结果在 batch 间稳定），但把”融合到什么程度”放在了不同位置。这也是阅读推理框架时一个可复用的观察角度：同一段逻辑放在哪一层，往往决定了框架的扩展方式。

下一篇的入口：从对齐表到分组 GEMM

本篇走到 moe_align_block_size 的输出为止：

1
输入：
2
  topk_ids      [T, k(+s)]    物理 expert id，padding 行为 -1
3
  block_size                   GEMM 的块粒度（Triton 实现通常是 64 或 128 行）
4
输出：
5
  sorted_token_ids     [max_num_tokens_padded]   按专家排序后的 (token*k + slot) 全局槽位号
6
  expert_ids           [max_num_m_blocks]        每个 block 负责的专家号
7
  num_tokens_post_pad  [1]                       总填充后的 slot 数
8
  cumsum（可选）        [E + 2]                   每桶起点

下一环是分组 GEMM 本身：fused_moe_kernel 如何用 expert_ids[block] 选出权重、用 sorted_token_ids 的 token*k+slot 反拆出 token 与 slot、在 block_size 行上做 A[token] × B[expert]，以及 decode 时”每专家只有十几行”的桶如何被 padding、如何与 BLOCK_SIZE_M 的取舍叠加。之后还有 EP 通信（all-to-all 与 DeepEP）与 EPLB 重排两篇——那时会回到本篇的”逻辑到物理映射”表，讲它如何被重新计算并下发到每一层。

参考资料

1. DeepSeek-V3 Technical Report
2. Mixtral of Experts
3. SGLang GitHub 仓库
4. vLLM 源码：fused_moe 路由与对齐实现