推理 INFRA 方法论（5）：MoE 推理：Batch、DeepEP、DeepGEMM 与负载均衡

博客Johney Zheng 的博客 · 全部文章索引 · 本文原文沿着 batch 这条主线分析专家覆盖率、通信量、tile 填充和负载倾斜并用端到端实测校验理论预算。Shape 与稀疏度的演变模型E(routed)k专家 inter稀疏度形态稠密 FFN (Llama-70B)——28672100%一个大 GEMMDeepSeek-V3 (2024)25681s20485.5%细粒度 sharedQwen3.5 (2026)512101s更小4.3%超细粒度 (shared 去了又回)DS-V4-Pro (2026)38461s20483.1% 原生 FP4趋势:E 涨 64×、单专家缩小 14×、稀疏度 25%→3%。后果:单专家 GEMM 从大而稠变小而碎——MoE 的一切系统问题(padding、EP、EPLB)都源于这个形态变化。MoE角度上看RoofLine下Batch/通信/算力的整体影响关系看Batch对计算效率的影响┌─→ AI ↑ → 吞吐 ↑(稠密模型 B* 之前近乎免费;MoE 还要乘 C(B),见下) │ ├─→ M_e ↑ → GEMM 从 padding 区推向饱和区 B ↑ ────────┤ ├─→ all2all 字节 ∝ B → 通信时间 ↑(跨机时成为新瓶颈) │ └─→ KV 池需求 ∝ B → 被显存卡住(EP 解锁) 实际运行点 min( B*, B_SLA, B_max_mem ) ↑吞吐最优 ↑TPOT上限 ↑显存上限B 是唯一同时出现在四条链上的变量。 所有 MoE 系统优化,本质都是在把 B 推高和消除推高 B 的副作用之间做工程——推高 B 靠 EP 解锁显存,副作用(通信涨、padding、不均衡)则各有对应的库来治。MoE 的“B* 之前免费”要打折一步内命中专家的比例C(B)1−(1−k/E)^B随 B 增长读取的权重也增长。令 W_routed 为路由专家参数数、W_shared 为每步必读的其他参数数则权重流量近似为[C(B)·W_routedW_shared]·bpp。只有 C(B) 接近 1 且尚未碰到计算等限制时才可能形成近似平坦区间。在全专家覆盖、忽略其他算子、相同分片方式的理想条件下令读权重时间与2·N_active·B/F相等得B*≈Ridge·W_total·bpp/(2·N_active)。这是权重 GEMM 的拐点估计不能直接等同服务吞吐最优 batch真实运行点还受 KV、attention、通信和 SLO 约束。链条恶化表现谁治机制all2all 字节 ∝ B跨机通信成为新瓶颈DeepEP把通信从加法项变成可被计算掩盖的项M_e 落在 padding 区为 pad 行付算力浪费 10×DeepGEMM让静态形状容纳动态负载pad 粒度降到真实 M_e专家负载不均all2all 同步屏障 → 最慢卡决定整层EPLB打散热专家把 max/mean 压回 1端到端实测校验理论下限与服务能力分开Qwen3.5-397B-A17B-FP88 卡、TP8input 1024/output 128关闭 radix cache静态显存比例 0.85decode graph 最大 batch 256。下表记录bench_one_batch_server的单批次采样结果。batchGraph decode tok/sTPOTeager TPOT11775.6 ms144 ms88968.9 ms141 ms32194816.4 ms146 ms64282822.6 ms143 ms128413630.9 ms—256562645.5 ms—严格 30ms TPOT 上限下B128 的 30.9ms 不达标已有采样中 B64 的 22.6ms 达标但并未定位最大达标吞吐。prefill 吞吐约 26K tok/sB1 TTFT 74ms。自然语言 prompt 的 MTP-3 接受长度 3.07–3.11B1/8/32 的 TPOT 为 2.45/5.1/7.6ms随机 token prompt 的接受长度 3.64 偏高不用于收益结论。未开 MTP 时 B≤8 已低于 10ms因此“10ms 必须 16 卡MTP”不成立。满权重读取下限397GB / 8 / 3.33TB/s 14.9ms假设本步读满全部权重小 batch 必须改用命中专家并集。B256 的 45.5ms 约为这个理想下限的 3.1 倍差值同时可能包含有效带宽不足、计算、通信等不能用相减推断“其他开销占 2/3”。DeepEP通信量的本质:dispatch B·C·d·1B(FP8) combine B·k·d·2B(BF16),其中C P[1−(1−1/P)^k]是去重后的目标卡数(EP32 时 7.18,不是 k8)。字节量随 B 线性涨,这一项无法消除,只能做三件事:① 减字节:dispatch 走 FP8(减半);同 rank 多专家共享一份 token 数据(去重 Ck) ② 压 α:IBGDA 让 GPU 的 SM 直接写网卡 WQE、敲 doorbell,绕过 CPU proxy —— 跨机 all2all 是高频小包,α 而非带宽才是主导项 ③ ★ 把加法变 max:hook 机制。send 发起后 RDMA 由网卡执行、不占任何 SM, 计算 kernel 照常排队,要数据时才调 recv hook ⟹ TBO(双 micro-batch 重叠)踩在这个机制上,层时间从 t_comm t_comp 变成 max(·,·)不占 SM只对跨机 RDMA 成立。机内走 NVLink 时数据由 SM 自己搬,实测 EP4(E256、k8、d7168)low-latency dispatch 32–50 μs、combine 30–76 μs,hook 重叠只藏住 20–27% 的通信时间。去重目标卡数实测 3.44–3.61,与公式 3.60 吻合。两套 kernel 对口两个阶段:normal(prefill,NVLink→RDMA 非对称域转发,吃带宽)/low-latency(decode,纯 RDMA 单跳 固定预分配 buffer → CUDA Graph 可 capture)。价值定位:DeepEP 不减少通信量,它让通信不再出现在关键路径上。所以它的收益上限 通信占步时间的比例,在跨机大 EP 场景可达 40-60%。DeepGEMMMoE decode 的核心矛盾是两个约束打架:约束 A(CUDA Graph):shape 必须静态 —— capture 时定死,replay 不能变 约束 B(MoE 现实): 每专家真实行数 masked_m[e] 每步都在变,且很小三种做法只有一种能同时满足:做法Graph 支持M_e5 时的浪费pad 到 M_max 的稠密 batched GEMM✅12.8× (算满 M_max)CUTLASS grouped (按真实 M)❌ shape 每步变12.8× (tile 太大)DeepGEMM masked✅3.2× (BM16) → 1.6× (BM8)这里是 FLOP 浪费核心一招:内存静态、计算动态。buffer 形状固定为[E_loc, M_max, K]满足 graph 要求而masked_m是一个 GPU tensor——地址固定、内容每步可变kernel 运行时读它决定实际算多少块m_size __ldg(masked_m[tile.e]); // 运行时读该专家真实行数 if (tile.m_blk * BM m_size) continue; // 纯 pad 块:一条 MMA 都不发配套两件事:两级累加解决 FP8 深 K 精度(TC 只连累 128 个 K,CUDA core 乘 scale 升 FP32;实测全正数据、K16K 时不升位误差约 25%,两级累加约 0.17%,等于 BF16 输出舍入);JIT 全特化可按形状选择 tile具体版本是否使用 expected_m 改变 BM 要实测本次 SM90 masked 路径改变 expected_m 并未改变实际 BM64。prefill 用 contiguous 布局和较大 tile 摊薄权重重读。为什么纯 pad 块不发 MMA还有浪费:跳过是按整个 tile 生效的,只要块里有一行真数据就要算满 BM 行,FLOP 浪费 ⌈M_e/BM⌉·BM / M_e。tile 还是访存瓶颈时,这部分浪费只是 FLOP,不占时间。BM 的临界值也是 Roofline:单个 tile 的算术强度 2·BM/bpp,与真实 M_e 无关。BM* Ridge×bpp/2以 FP8 的参考参数计算约为 37代入实测有效吞吐后约为 42——BM 低于它时 tile 仍是访存瓶颈,pad 行藏在权重读取里近乎免费;高于它后 tile 翻成算力瓶颈,pad 行开始真金白银烧时间。实测(sglang Triton fused_moe,W8A8 block-128,8 专家 gate_updown 共 352 MB,M_e ≤ 16):BM16/32/64/128 分别 118/159/202/445 μs,对应 3.0/2.2/1.7/0.8 TB/s,BM 128→16 快 3.8×;实际 kernel 在 BM≈16–24 就开始掉速,比理论拐点早。DeepGEMM 0.2.2 的 SM90 masked 路径实际用 BM64,小 M_e 下只有 1.8 TB/s。★ 两库咬合:DeepEP 的 recv_count[num_local_experts] 就是 DeepGEMM 的 masked_m,接收 buffer 直接作 GEMM 输入,零拷贝零重排。EPLBall2all 是同步屏障 → MoE 层时间 最忙卡的时间 → 慢S₀ max_load/mean_load倍。倾斜有两个独立来源,只有一个能治:① 统计涨落:S₀ ≈ 1 √(2lnP) · CV_gpu,其中 CV_gpu CV_e/√e(卡内 e 个专家自平均) B_local48 时地板约 ±13% —— ★ EPLB 治不了,只能靠加大 batch ② 路由热点:S₀ ≈ 热冷比 r(坏 placement 时) ⟹ EPLB 复制热专家 LPT 装箱打散,把 S₀ 压回 ~1.1收益换算(注意两级摊薄):MoE 层提速 S₀/S₁ 端到端 1 / ( f/(S₀/S₁) (1−f) ) f 可均衡占比 ↑ EPLB 同时均衡 GEMM 和每卡 a2a 接收量(都 ∝ 专家负载) 治不了的:总字节量、α 地板、attention场景e(每卡专家)CV_eS₀/S₁f端到端EP32,B48,跨机8~0.351.180.61.10×SGLang EP72 decode4≥1.3~2.60.852.54×★ EPLB 的收益由三个因子决定,缺一不可:热点程度(CV_e)、每卡专家数(e 越小越无自平均)、可均衡占比(f)。e8 大 batch 均匀路由 10% 就是天花板,不是实现问题;要看到 2× 级收益,需要 EP64/e≤4 真实重尾路由 配冗余专家。系列导航章节文章编排日期0Roofline用算力与带宽建立性能预算2026-01-011通信模型从集合原语到跨机拓扑2026-02-132分布式推理TP、DP、EP、CP 与 PP 的取舍2026-03-293推理硬件GPU、NPU 与存储层级2026-05-114Attention 演变从 GQA、MLA 到混合状态2026-06-245MoE 推理Batch、DeepEP、DeepGEMM 与负载均衡本篇2026-08-066PD 与 EPD阶段分离、状态传输与资源配比2026-09-197核心优化CUDA Graph、投机推理、Overlap 与量化2026-11-01上一篇Attention 演变从 GQA、MLA 到混合状态 · 下一篇PD 与 EPD阶段分离、状态传输与资源配比