:INT8 高吞吐/低延迟 EP 代码梳理、疑问、流程差异、DeepEP 通信差异)
★ FEATURED ARTICLE 目录1 vllm/distributed/device_communicators/all2all.py2 vllm/model_executor/layers/fused_moe/all2all_utils.py3 vllm/model_executor/layers/fused_moe/batched_deep_gemm_moe.py3.1 增加 N、K 参数 放宽 init assert3.1.1 问题这里为什么只需要更新N3.2 替换算子名字4vllm/model_executor/layers/fused_moe/config.py5 vllm/model_executor/layers/fused_moe/deep_gemm_moe.py5.1 初始化函数5.2 permute:token重排两段代码是在 MoE 第一次 GEMM 之前为 permute按 expert 重排 token 准备 buffer 大小和形状。6 vllm/model_executor/layers/fused_moe/deep_gemm_utils.py7 vllm/model_executor/layers/fused_moe/deepep_ht_prepare_finalize.py8 疑问高吞吐和低延迟的流程差异问题9 疑问scatter和gather具体做了什么inv_perm和topk_ids是什么9.1 先建立画面dispatch 之后本卡有什么9.2 topk_ids 是什么9.3 inv_perm 是什么名字什么意思9.4 scatter / permute 到底在干什么M 和 M_sum 啥关系9.4.1 为什么要重排9.4.2 具体怎么排小例子先忽略 256 对齐9.4.3 M_sum 是什么不是 M_num9.5 scatter 之后 GEMM再 gather9.6 gather 再讲一遍unpermute 加权9.7. 一张总流程图9.8. 几个容易混的点10 疑问:M_sum_upper (M * num_topk) local_num_experts * (alignment - 1)为什么要加上local_num_experts * (alignment - 1)11 高吞吐和低延迟deepep代码有什么区别11.1 为何 Prefill → HTDecode → LL11.2 布局与 GEMMStandard/contiguous vs Batched/masked11.3 通信 BufferHT 固定中转 紧凑区 vs LL 大 RDMA/XDP11.3.1 低延迟按容量预开的大 RDMA11.3.2 高吞吐较小固定中转 按真实 N 开的紧凑区11.4.1 HTnotify_dispatch 在真正发激活之前11.4.2 LL没有同款前置 notify_dispatch11.5. HT 两趟搬数LL 无单独 permute11.5.1 HT 数据落点逻辑上两趟11.5.2 LL11.5.3 加权发生在哪和通信相关的收尾11.6. 256 对齐、量化位置差异11.6.1 256 对齐11.6.2 量化时机以常见路径为例11.7. 对照表收尾abstract:只有高吞吐要用permute函数重新排列tokenpermute英/ pəˈmjuːt / 美/ pərˈmjʊt /改变……的次序重新排列scatter 英/ ˈskætə(r) / 美/ ˈskætər / 撒播撒使散开使散布在各处物理散射电磁辐射或粒子棒球被击中但没有得分的有效投球HTdispatch 和 permute 分工量化 → dispatch(跨卡) → permute(本卡 scatter) → GEMM×2 → unpermute(本卡 gather加权) → combine(跨卡发回)这个流程刚看完很量化dispatch, scatter, ;gemm, gather combineLLdispatch 已带 batched layoutlow_latency_dispatch (跨卡 按 expert 摆好 (E,T,H)) → masked GEMM×2 → low_latency_combine (跨卡 加权归约)LL 里 permute 的职责被 dispatch 的布局 masked GEMM 的expert_num_tokens吸收了unpermute 的职责被 combine 吸收了。高吞吐和低延迟量化有什么区别:高吞吐外部量化低延迟是deepep内部量化256干什么用的256就是传进去给对齐用的然后我们的deepep就是要求按照256对齐并且只有高吞吐要求256对齐低延迟不是256对齐的。之前的一个bug是因为什么我现在好像理解了就是这个256传进来其实就是对齐用的但是还有个问题就是我们的高吞吐是deepep外部也就是我们的vllm量化所以这个错误其实就是因为我们本来传进来blockshape为了对齐用结果量化的时候判断了blockshape不是none那么就报错了对吧所以修改应该是可以在高吞吐里面写死比如定义个align_shape[256, 256]一类的然后不要用blockshape这样就不会混淆了对吧m_indices就是 expert_ids:一维形状(M_sum,)长度等于aq_out的行数一行对应一个专家号或-1。含义:expert_ids就是每行激活属于哪个专家ep_scatter把 token 按 expert 排成连续布局后用这份表告诉 grouped GEMM「每一行归谁」。[ 0,0,0,...,0, -1,-1,...,-1 | 1,1,1,...,1, -1,-1,...,-1 | 2,... ]← expert0 真实行 → ← pad到256 → ← expert1 → ← pad →每个 expert 一块连续行不够 256 就用-1补齐DeepGEMM 只算有专家号的行遇到-1跳过。expert_ids[i] e → 给 expert e 算expert_ids[i] -1 → padding / 无效DeepGEMM 跳过inv_perm inverse permutation逆置换 / 反向索引。scatter 把 token 按 expert 排到连续布局inv_perm记下「原来每个(token, topk)被放到了哪一行」后面 gather / reduce 再按这个索引捞回来。inv_perminverse permutation就是表示某个token的某个专家在整个结果的哪一行和topk_ids同形状一般是 2 维(num_tokens, topk)例如inv_perm[t, k] dest_row第t个 token 的第k个 expert 选择scatter 到布局的第dest_row行。1 vllm/distributed/device_communicators/all2all.py这个文件其实修改了deepep所用的sm数量下面的是高吞吐对应的RDMA的字节数等这些内容这个是所有的什么int8 fp8 w4a8共用的。改动说明num_sms基类20 → 30DeepEP 通信用多少 SMnum_nvl_bytes节点内 NVLink buffer 加大约 1GBnum_rdma_bytes跨节点 RDMA buffer 加大约 0.5GB仅 internodenum_qps_per_rankRDMA 通道数Queue Pair10 → 30intranode 时num_sms进一步设为 60与 dtype 无关INT8、FP8、W4A8 等只要走 DeepEP EP都会用到这层通信不是 INT8 专用。不全是「所有 backend 同一套改法」基类num_sms30对 HT/LL 都生效buffer / RDMA / QP 那一大段主要在DeepEPHTAll2AllManager高吞吐里。低延迟 LL 有自己的 buffer 初始化逻辑。2 vllm/model_executor/layers/fused_moe/all2all_utils.py它是 EP 通信的「工厂 / 路由」根据当前 MoE 配置的 backend创建对应的PrepareAndFinalize对象。高吞吐和低延迟都会走这里入口是同一个函数MoE 层初始化 modular kernel → quant_method.maybe_make_prepare_finalize() → all2all_utils.maybe_make_prepare_finalize() → 根据 backend 创建对应的 PrepareAndFinalize 对象 ├── deepep_ht → DeepEPHTPrepareAndFinalize └── deepep_ll → DeepEPLLPrepareAndFinalize然后这里修改的是放宽 FP8 条件只要是 FP8 dtype 就use_fp8_dispatchTrue适配 channelwise [256,256]block新增 INT8use_int8_dispatch (quant_dtype torch.int8)并传给 LL 的 Prepare/Finalize。为什么ht的时候不需要传这种inT8 fp8参数3 vllm/model_executor/layers/fused_moe/batched_deep_gemm_moe.py3.1 增加 N、K 参数 放宽 init assertEP 下权重经weight8bit_nt_kpack2_marlin2重排后w1.size()的物理 shape 与 GEMM 逻辑维度不一致。因此在compressed_tensors_moe_marlin中保存self.N、self.K经select_gemm_impl传入 Expert在apply()里覆盖。同时放宽 upstream 仅支持 FP8G128 的__init__assert以支持 INT8 与block_shape[256,256]。3.1.1 问题这里为什么只需要更新N因为这里的K其实是激活的维度他来自_, N, K w1.size() # ① 从 weight可能错 E, max_num_tokens, N, K, _ self.moe_problem_size( hidden_states, w1, w2, topk_ids) # ② 重算 # moe_problem_size 内部 # N w1.shape[1] ← 仍可能错 # K hidden_states.size(-1) ← 从 (E,T,H) 的 H 来一定对 if self.N 0: N self.N # ③ 只修 N所以这里的K没问题只需要修改N就行了但是对于后面的高吞吐的他其实NK都需要因为高吞吐那里没有这个self.moe_problem_size函数。LLHTdispatch 后 activation(E, T, H)(M, H)→ permute 成(M_sum, K)N 从哪来w1.shape[1]→ 可能错w1.shape[1]→ 可能错K 从哪来hidden_states.size(-1)✅w1.shape[2]❌ 重排后错要不要self.N✅✅要不要self.K❌ 已有正确 K✅ 没有别的来源3.2 替换算子名字这里就是将社区的算子名称替换为我们公司自己的算子了没什么4vllm/model_executor/layers/fused_moe/config.py这个是当初256的问题先不用管。5 vllm/model_executor/layers/fused_moe/deep_gemm_moe.py这个文件就是高吞吐所对应的文件5.1 初始化函数这块就是放宽了关于int8的量化类型的判断还有传入了N K ,后面FP8的高吞吐这里就是需要修改的。5.2 permute:token重排两段代码是在 MoE 第一次 GEMM 之前为 permute按 expert 重排 token 准备 buffer 大小和形状。这个 函数是在做什么DeepGEMM HT 路径里token 先按 router 分到各 expert再 按 expert 聚成连续块 做 grouped GEMM。281–291 就是在算这次 forward 需要多少行M_sum以及 从预分配 workspace 里切出哪一块当 permute 输出。原始 token: [t0, t1, t2, ...] (M 个 token × topk 路由) ↓ permute 按 expert 排: [expert0 的 token... | expert1 的 token... | ...] (M_sum 行) ↓ grouped GEMM (w1)6 vllm/model_executor/layers/fused_moe/deep_gemm_utils.py这个文件主要就是deep_gemm_utils.py HT 专用的 permute / unpermute 工具compute_aligned_M算M_sumdeepgemm_moe_permute按 expert scatter 重排deepgemm_unpermute_and_reducegather 还原 topk 加权求和具体细节先不展开了7 vllm/model_executor/layers/fused_moe/deepep_ht_prepare_finalize.py8 疑问高吞吐和低延迟的流程差异问题只有高吞吐要用permute函数重新排列token英/ pəˈmjuːt /美/ pərˈmjʊt /HTdispatch 和 permute 分工量化 → dispatch(跨卡) → permute(本卡 scatter) → GEMM×2 → unpermute(本卡 gather加权) → combine(跨卡发回)LLdispatch 已带 batched layoutlow_latency_dispatch (跨卡 按 expert 摆好 (E,T,H)) → masked GEMM×2 → low_latency_combine (跨卡 加权归约)LL 里 permute 的职责被 dispatch 的布局 masked GEMM 的expert_num_tokens吸收了unpermute 的职责被 combine 吸收了。高吞吐和低延迟量化有什么区别:高吞吐外部量化低延迟是deepep内部量化256干什么用的256就是传进去给对齐用的然后我们的deepep就是要求按照256对齐并且只有高吞吐要求256对齐低延迟不是256对齐的。之前的一个bug是因为什么我现在好像理解了就是这个256传进来其实就是对齐用的但是还有个问题就是我们的高吞吐是deepep外部也就是我们的vllm量化所以这个错误其实就是因为我们本来传进来blockshape为了对齐用结果量化的时候判断了blockshape不是none那么就报错了对吧所以修改应该是可以在高吞吐里面写死比如定义个align_shape[256, 256]一类的然后不要用blockshape这样就不会混淆了对吧9 疑问scatter和gather具体做了什么inv_perm和topk_ids是什么inv_perm是什么意思什么是inv_perm这个是两个单词的缩写吗inv_perm[t, k]这个inv_perm是个二维的吗就是比如对于第t个token它对应的肯定是多个expert吗所以应该是二维的吧然后这个列表的那个值就是说在第几行那么topk_ids也能给我解释一下吗是不是也是二维的里面是个什么东西还有原来的token是M行我知道就是M个token然后每个token一行这个我懂这就是dispatch刚发过来的这里的顺序其实就是token的先后顺序然后你说的permute或者scatter我还是有点不太懂包括这个计算的M_sum行这个能再给我解释一下吗难道是将原来是先后顺序的token现在按照专家排列吗具体怎么排列的这个M_num又是什么意思gather 阶段具体做什么对每个原始 token t共 M 个看 router 选的 topk 个 experttopk_ids[t, 0..topk-1]对每个有效 expert用 inv_perm[t, k] 查到这个 (token, expert) 在 重排缓冲区 里是第几行从 mm2_out 那一行 读出 hidden vector乘以 topk_weights[t, k]把 topk 路结果 加起来写入 output[t, :]形状 (M, K)K hidden_size也就是按原 token 顺序收集 加权合并MoE 里常说的 unpermute reduce。Triton kernel _fwd_kernel_ep_gather 里核心逻辑就是还有这个gather的过程也是重新给我讲一下用一个具体小例子把topk_ids、inv_perm、scatter 的M_sum行布局以及 gather 怎么还原串起来说明。9.1 先建立画面dispatch 之后本卡有什么DeepEP dispatch 之后本 rank 拿到张量形状含义aq/ hidden(M, K)M 个 token每行一个 hidden 向量topk_ids(M, topk)每个 token 选了哪 topk 个 experttopk_weights(M, topk)每个 expert 的 router 权重加起来≈1你说的对dispatch 刚过来时就是 M 行按本卡收到的 token 顺序排第 0 行、第 1 行……。此时还没有按 expert 分组只是「M 个 token 各带一条 routing 信息」。9.2topk_ids是什么是二维的形状(M, topk)。第 1 维t第几个 token0 ~ M-1第 2 维k该 token 的第几个 expert 选择0 ~ topk-1topk_ids[t, k]里存的是 expert 编号global expert idEP 下会 remap这里先当「专家 ID」理解。例如 MoE 每个 token 选 topk2 个 experttopk_ids[0, :] [专家3, 专家7] ← token0 要去专家3和专家7topk_ids[1, :] [专家3, 专家1] ← token1 要去专家3和专家1topk_ids[2, :] [专家7, 专家1] ← token2 ...一个 token 对应多列topk 列因为 MoE 是「多个 expert 的输出加权相加」不是只选一个。topk_weights[t, k]与之一一对应表示「token t 走 experttopk_ids[t,k]时乘多少权重」。9.3inv_perm是什么名字什么意思是两个词inv inverse逆perm permutation置换 / 重排合起来逆置换表——scatter 把 token 重排到别的行号后gather 需要查表找回「原来 (token t, 第 k 个 expert) 被放到了第几行」。也是二维的形状和topk_ids一样(M, topk)。inv_perm[t, k] 一个整数行号含义tokent的第k个 expert 路由在 重排缓冲区scatter 输出 / GEMM 输入输出里是 第几行。不是 expert 编号不是 token 编号就是 行 index0, 1, 2, …, M_sum-1scatter 写入代码里叫output_indexdeep_gemm_utils.pyLines 173-176tl.store(output_index token_id * output_index_stride0 topk_index,dest_token_index_int32,)也就是每处理一次 (token t, slot k)往inv_perm[t,k]写「拷贝到了第几行」9.4 scatter / permute 到底在干什么M 和 M_sum 啥关系9.4.1 为什么要重排Grouped GEMM 希望数据是同一个 expert 的 token 连续放在一起一次 kernel 里按 expert 分组算矩阵乘。dispatch 之后是 按 token 顺序 的 M 行不是按 expert 分组。scatter 按 expert 把 M 行「拆开、拷贝、拼成新顺序」。可以记从「按 token 排」变成「按 expert 排」。9.4.2 具体怎么排小例子先忽略 256 对齐设M3 个 tokentopk2本卡只有 3 个 local expertE0, E1, E2tokentopk_ids[t,0]topk_ids[t,1]t0E1E2t1E0E1t2E1E2每个 (token, expert) 要在缓冲区里占 1 行拷贝一份 hidden。一共 M × topk 3×2 6 份但按 expert 分段 摆【E0 段】行0: token1 的 hidden因为 t1 选了 E0【E1 段】行1: token0 的 hiddent0 选 E1行2: token1 的 hiddent1 选 E1行3: token2 的 hiddent2 选 E1【E2 段】行4: token0 的 hiddent0 选 E2行5: token2 的 hiddent2 选 E2重排缓冲区一共 6 行 → 这个例子里 M_sum 6无 padding 时。对应的inv_perminv_perm[0,0]1 inv_perm[0,1]4 # token0 → E1在行1, E2在行4inv_perm[1,0]0 inv_perm[1,1]2 # token1 → E0在行0, E1在行2inv_perm[2,0]3 inv_perm[2,1]5 # token2 → E1在行3, E2在行5同一条 token 的 hidden 可能被拷贝多行选了几个 expert 就拷贝几次每行对应一次 expert 计算。9.4.3 M_sum 是什么不是 M_num代码里叫M_sumsum 的意思不是 M_num。M_sum scatter 之后重排缓冲区的总行数用于aq_out/mm2_out的第 0 维grouped GEMM 的「逻辑 batch 行数」怎么算每个 expert 先数有多少个 (token, slot) 指向它每个 expert 的行数 向上对齐到 256DeepGEMM / HT 要求把所有 expert 的对齐后行数加起来 M_sum所以通常M_sum ≥ M × topk多出来的行是 padding对齐用不一定有真实 token。M原始 token 个数dispatch 后的(M, K)M_sum按 expert 重排 对齐后的行数(M_sum, K)9.5 scatter 之后 GEMM再 gatherscatter 后aq_out / mm2_out : (M_sum, K) ← 按 expert 分组 paddingtopk_ids : (M, topk) ← 仍是「每个 token 选了谁」未变inv_perm : (M, topk) ← scatter 填好的「行号对照表」中间 GEMM1 → SiLU → GEMM2在 M_sum 行 上算得到mm2_out (M_sum, K)每一行 「某个 token 某条 expert 路由」算完 expert 后的 hidden。9.6 gather 再讲一遍unpermute 加权目标从(M_sum, K)按 expert 排的结果变回(M, K)按 token 排并完成 MoE 的 加权求和。对 每个原始 tokent共 M 个output[t, :] 0对 k 0, 1, ..., topk-1:expert topk_ids[t, k]若 expert 在本 rank 有效:row inv_perm[t, k] ← 在 mm2_out 的第几行w topk_weights[t, k] ← router 权重output[t, :] w * mm2_out[row, :]逐步对应步骤做什么1看topk_ids[t, :]token t 用了哪几个 expert2查inv_perm[t, k]该路由在mm2_out的第几行3读mm2_out[row, :]4乘topk_weights[t, k]并累加5得到output[t, :]形状(M, K)这就是 gather收回来 reduce加权求和。9.7. 一张总流程图dispatch 后:hidden (M, K) 按 token 顺序第 t 行 token ttopk_ids (M,topk) 每行这个 token 去哪些 expert│ scatter / permute▼aq_out (M_sum, K) 按 expert 分段同 token 可能占多行inv_perm (M,topk) inv_perm[t,k]「t 的第 k 路 expert 在第几行」│ GEMM1 → act → GEMM2▼mm2_out (M_sum, K) 每行 某 (token, expert) 的 expert 输出│ gather / unpermute▼output (M, K) 又回到按 token 顺序每行 topk 路加权求和9.8. 几个容易混的点问题答案inv_perm是二维吗是(M, topk)和topk_ids同形状inv_perm[t,k]是什么行号不是 expert idtopk_ids[t,k]是什么expert 编号一个 token 几列topk 列选几个 expert 就有几列M 和 M_sumMtoken 数M_sum重排后总行数≥ M×topk含对齐 paddinggather 在干什么用inv_perm从mm2_out按 token 收回来用topk_weights加权相加10 疑问:M_sum_upper (M * num_topk) local_num_experts * (alignment - 1)为什么要加上local_num_experts * (alignment - 1)就是因为对齐的时候是每个专家根据256对齐而不是所有的行数总和根据256对齐所以这里就需要加上个local_num_experts * (alignment - 1)11 高吞吐和低延迟deepep代码有什么区别11.1 为何 Prefill → HTDecode → LLPrefillDecodetoken 量大小常接近并发请求数优化目标tokens/s吞吐单步延迟 / TPOTshape可变通常可接受最好固定CUDA GraphHThigh_throughput 适合 Prefill用精确 count 做紧凑布局打满带宽和大 M 的 contiguous GEMM一次 CPU sync 相对整步可忽略。LLlow_latency 适合 Decode固定(E,T,H)、少 CPU sync、好上 Graph小 batch 时若再走 HT 的 256 对齐 变长整理固定开销会占主导。11.2 布局与 GEMMStandard/contiguous vs Batched/masked两边 dispatch 之前激活都是 Standardx: (M, H)topk_ids / topk_weights: (M, topk)差别在 dispatch 之后HTLLdispatch 后激活紧凑 2D(N_recv, H)Standard 系batched 3D(E, T, H)框架 permute需要本卡 scatter 成专家连续不需要dispatch 已按 expert 摆好GEMMcontiguous grouped GEMM m_indicesmasked grouped GEMM expert_num_tokens/masked_munpermute本卡 gather 加权职责更多在low_latency_combine框架侧流程对照HT:量化(视路径) → HT dispatch → permute/scatter → GEMM×2→ unpermute/gather(加权) → combine(跨卡求和发回)LL:low_latency_dispatch(跨卡 收成 (E,T,H)常内核内量化)→ masked GEMM×2→ low_latency_combine(跨卡 加权归约)(E,T,H)是 LL dispatch 的接收结果不是发送前就已经 batched。11.3 通信 BufferHT 固定中转 紧凑区 vs LL 大 RDMA/XDP两边都有deep_ep.Buffer但 怎么开、开多大完全不同。11.3.1 低延迟按容量预开的大 RDMAnum_rdma_bytes ≈ get_low_latency_rdma_size_hint(max_tokens_per_rank, hidden, num_ranks, num_experts)low_latency_mode True大小跟max_num_batched_tokens/ max tokens 强相关常落在可 RDMA 的显存例如环境里的 XDP ~4GB过大就报错提示减小--max-num-batched-tokens或改用 HT接收布局按专家固定槽(E, ranks * max_tokens, H)就地给 masked GEMM 用11.3.2 高吞吐较小固定中转 按真实 N 开的紧凑区low_latency_mode Falsenum_nvl_bytes 配置固定如 VLLM_DEEPEP_BUFFER_SIZE_MBnum_rdma_bytes 固定较小 / 节点内甚至可为 0不会用 LL 那套 hint 按 max token 把 RDMA 撑到数 GB。HT 是两块不同内存中转 Buffer紧凑recv_x/expert_x寿命Buffer 长期持有每次 dispatch 按N_recv临时empty类型预注册通信显存普通 CUDA tensor不吃那套 4G XDP 公式作用多路队列边收边腾收齐后的紧凑激活中转队列 有限不能一次性塞下整个 Prefill靠 多 rank 并发 流水 drain 搬进紧凑区。11.4. notifyHT 前置换 count vs LL kernel 内 flag/count这里最容易混prof 里 LL 也能看到 notify但 不是 HT 那种两阶段 API。11.4.1 HTnotify_dispatch在真正发激活之前notify_dispatch 交换「会收多少」 → CPU sync 拿到精确 num_recv_tokens / 每专家 count → torch::empty({N_recv, H}) 开紧凑区 → 再发真正的 token 数据目的为变长紧凑 buffer 申请服务。这也是 Prefill 可接受、Decode 很疼的那次 host sync。11.4.2 LL没有同款前置notify_dispatch调用侧基本直接low_latency_dispatch(...)。槽位已按上限预开不需要「先问清精确 N 再 malloc 紧凑区」。但 dispatch kernel 内部仍有 notify/flag往对端固定槽 put 数据 → atomic/flag 通知 count数据到了、到了多少 → 接收端等 flag 后再读、累加 packed_recv_countHTnotify_dispatchLL 内部 notify形态独立前置阶段嵌在 LL dispatch kernel 里目的换精确个数开紧凑区到达/完成信号GPU 侧CPU sync 开 buffer要不要能否去掉去掉则无法按真实 N 开区去掉会读脏数据 / 不知有效行数11.5. HT 两趟搬数LL 无单独 permute11.5.1 HT 数据落点逻辑上两趟① 固定中转 Buffer通信序 / 多路队列 ↓ DeepEP 收包/打包 kernel不是框架 cudaMemcpy ② 紧凑 expert_x (N_recv, H) ← 通常按源 rank 等通信序尚未按专家聚齐 ↓ 框架 ep_scatter / permuteTriton 或 LightOP scatter kernel ③ 按专家连续 workspace可含 256 paddingm_indices / inv_perm ↓ contiguous DeepGEMM要点紧凑 ≠ 已 permute紧凑只是通信收完的 2D 落点第二次不是裸memcpy是 按 routing scatter再后面才是 GEMM → unpermute → combine11.5.2 LL(M,H) → low_latency_dispatch → (E,T,H) → masked GEMM → low_latency_combinepermute / unpermute 的职责被 dispatch 布局 combine 吸收没有 HT 那套「紧凑 → 再 scatter」的框架步骤。11.5.3 加权发生在哪和通信相关的收尾乘topk_weights跨 rank 加总HT本卡 gather/unpermutecombine常topk_weightsNone加 partialLL多在low_latency_combine同在 combinetoken 选了不同 rank 的专家时各卡先形成加权 partialcombine 再跨卡相加——HT 是「先乘后跨卡加」LL 更常揉进 combine。11.6. 256 对齐、量化位置差异11.6.1 256 对齐只服务 HT contiguous每个专家段 pad 到 256 倍数写入m_indices-1为 paddingLL 不是 256 对齐那套LL 的第二维是固定max_tokens容量不要和量化block_shape混用(26) 里那个 bug 的根因11.6.2 量化时机以常见路径为例常见做法HT INT8per-tokenHCU框架先量化再 HT dispatchHT 部分 channel FP8可能 先 dispatch再框架量化LL常在low_latency_dispatch内核内quant_type所以「HT 一定外部、LL 一定内部」大体对但 HT 内部仍有分支写 INT8 HT 时用「量化 → dispatch → scatter → …」没问题。11.7. 对照表收尾维度HT高吞吐 / PrefillLL低延迟 / Decode发送前激活(M,H)Standard(M,H)Standarddispatch 后紧凑(N_recv,H)batched(E,T,H)通信显存固定较小中转 临时紧凑区大块 RDMA/XDP按 max token精确 count→CPU要开紧凑区不要count 留 GPUnotify前置notify_dispatchkernel 内 flag/count框架 permute要不要GEMMcontiguous m_indicesmasked masked_m256 对齐要不要加权多在本卡 gather多在 combineCUDA Graph难变长易固定 shape
阅读完成 · 觉得有帮助？
