完整交互版https://besproma.github.io/deepinfer-ascend-analysis/代码快照774ce24分析日期2026-10-21分析范围并行调度、模型级优化、昇腾适配、芯片内数据流、源码证据与公开性能数据DeepInfer 是 DeepSeek 面向 DeepSeek-V4.1-Flash 深度定制的高性能推理框架。它与通用推理框架的主要差别不是又增加了一套模型接入接口而是把模型结构、稀疏 Attention、MoE 路由、KV 规则和服务调度放在同一个设计闭环中并分别实现 NVIDIA 与昇腾后端。本文重点回答三个问题DeepInfer 为什么要做这些优化昇腾适配到底落在哪些工程层以及公开性能数据能够说明什么。一、为什么要重新组织推理路径一次大模型请求看似只是“输入 Prompt、输出 Token”实际会经过一条很长的关键路径Prefill 阶段处理整段 Prompt并产生首 Token。MoE 层根据路由把 Token 分发到不同专家再汇合结果。Decode 阶段逐个 Token 自回归生成。历史上下文在 KV Cache 中复用缓存未命中时还要恢复。CPU 侧持续完成请求检查、批处理拼装和元数据准备。这条链路的耗时往往不只来自矩阵计算。长 Prompt 可能只落在一组设备上MoE 的热点专家可能让最忙 rank 决定整轮时间Decode 每一步都很小却要重复读取权重和启动 kernelCPU 排批、集合通信和中间张量写回还会制造额外等待。DeepInfer 的核心思路是让模型已经知道的形状与规则进入执行系统提前消除这些等待。关键路径常见问题DeepInfer 对应机制目标长上下文 Prefill单条请求难以拆到更多设备Balanced CP缩短首 Token 关键路径MoE 专家路由热点专家与跨 rank 分发形成等待EP、DeepEP、EPLB扩大专家池并压低通信与尾部负载逐 Token Decode串行步数多小 batch 算力利用率低DSpark一次验证多个候选 TokenKV 缓存恢复已失效历史仍被完整重放SWA bounded replay只处理当前输出需要的最短后缀CPU / NPU 协同设备执行期间 CPU 在准备下一批异步批调度让排批与设备执行重叠MoE 数据路径多个 kernel 反复写回和读取中间结果MegaMoE把分发、GEMM、SwiGLU、汇合接成连续数据流这些机制并非每一项都由 DeepInfer 首次提出。它的差异化在于围绕一个固定模型把并行方式、通信布局、缓存语义与硬件数据流一起调整。二、七项关键技术如何减少等待1. Balanced CP让一条长请求也能多卡并行传统的数据并行主要增加同时处理的请求数量却不会自动拆开单条长 Prompt。当一条请求的上下文很长时它仍然可能由一组设备独自承担某个 rank 分到更多 Token其他 rank 就要等待形成木桶效应。Balanced CP 按上下文切分一条请求并尽量平衡各 rank 的 Token 数和计算量。它扩大单请求可用的计算资源同时缩小不同 rank 的完成时间差因此优化的是首 Token 阶段的尾延迟。2. MoE EP 与 DeepEP把专家容量扩到集群MoE 模型包含大量专家。专家并行把不同专家分布到不同 rank从而扩大可容纳的权重和可用算力。但 Token 必须先通过 dispatch 送到目标专家计算后再通过 combine 返回原 rank通信路径处理不好新增算力会被跨卡等待抵消。DeepEP-Ascend 负责昇腾上的分发与汇合并接入 HCCL、HCOMM、UBMEM 和 URMA 等通信能力。它的价值来自两件事同时成立专家计算扩展到更多设备跨设备通信成本得到控制。公开验证覆盖 EP128。这里的 EP128 表示 128 个专家并行 rank不是 128 个专家。公开数据中dispatch 峰值约为 313–320 GB/s。3. EPLB让专家布局跟随真实热点专家铺到更多设备之后总量通常不再是主要矛盾局部热点才是。某些专家承接的 Token 明显更多对应 rank 就会成为新的长尾。EPLB 先观察真实 Token 的负载再计算新的专家布局并通过重排压缩最忙 rank 的排队时间。当前实现已经覆盖负载统计、方案求解和周期重排在线权重迁移保留平台对称内存接口便于继续利用昇腾的统一总线能力。4. DSpark把一步一 Token 改成一次核验多个标准自回归 Decode 每轮只前进一个 Token。生成序列较长时串行轮次很多小 batch 下设备算力难以铺满但每轮仍要承担权重读取、kernel 启动和调度成本。DSpark 先生成多个轻量候选 Token再由主模型一次核验。主模型接受正确前缀后一次迭代可以前进多个 Token。实际收益取决于候选接受率、核验成本和额外调度开销。昇腾算子目录中已经包含 DFlash、采样和相关调度路径说明该机制已进入硬件后端实现而不只停留在模型语义层。5. SWA bounded replay不要重算已经失效的历史滑动窗口注意力只关注最近的上下文窗口。缓存恢复时如果仍按通用路径完整重放历史窗口之外的 Token 虽然已经不影响当前输出却还会消耗 KV 读取、数据搬运和 Prefill 计算。SWA bounded replay 只重放覆盖当前注意窗口和最大回溯长度的必要后缀。V4.1 默认保留 140 个 Token具体收益取决于窗口配置、缓存复用率和命中率。6. 异步批调度下一批在上一批执行时准备好如果 CPU 和 NPU 严格轮流工作每一轮都要先完成请求检查、batch 拼装和元数据准备再交给设备执行。Decode 单批执行时间较短时这部分 CPU 等待会占据明显比例。异步调度让 CPU 在 NPU 执行当前 batch 时准备下一批。理想情况下一轮耗时更接近“CPU 准备时间”和“NPU 执行时间”中较长的一项而不是两者相加。NPU Graph 还能通过捕获与 replay 降低 kernel 启动开销。7. MegaMoE把多次 kernel 往返接成一条数据流dispatch、grouped GEMM、SwiGLU 和 combine 如果分别作为独立 kernel 执行每一步都要把中间结果写回再由下一步读取。这会增加访存、缓冲区占用、kernel 启动和同步点。MegaMoE 让通信布局、专家矩阵形状和中间结果传递彼此匹配减少中间张量的往返并创造通信与计算重叠的机会。它优化的是整条 MoE 数据路径而不是单个 GEMM 的峰值。三、DeepInfer 与 vLLM、SGLang 的关系DeepInfer 最初主要基于 vLLM 构建同时吸收了 SGLang、TensorRT-LLM 等框架的设计思路。它并非从零重写所有通用能力也不是替代所有推理框架。与 vLLM、SGLang 的主要差异可以概括为通用框架优先覆盖更多模型、硬件和接口DeepInfer 则只围绕 DeepSeek-V4.1-Flash 做单模型垂直优化。固定模型带来了更强的协同设计空间算子可以围绕固定 head 形状、稀疏模式和精度格式做极致调优。通信可以理解 MoE 路由和 GEMM 时序而不只提供通用 all-to-all。调度可以把 Draft Token、SWA 重放、缓存布局和 NPU Graph 一起考虑。昇腾后端可以针对 CUBE、VECTOR、L1、UB 和 FixPipe 重新组织数据流。代价同样明确模型泛化能力较弱版本耦合更强每代硬件都需要维护独立实现。设计维度通用推理框架常见策略DeepInfer 策略技术含义模型范围多模型、多模态、广泛兼容当前聚焦 DeepSeek-V4.1-Flash可以做模型级协同跨模型复用需重新适配Attention通用 PagedAttention / FlashAttentionDSA 稀疏索引、FP8 / FP4 KV、固定 head 形状稀疏场景优化空间更大非目标模型不适用MoE 通信通用 all-to-all 与集合通信DeepEP 与 MegaMoE面向专家 GEMM 的通信和计算协同执行调度通用连续批处理异步调度、DSpark、SWA 有界重放针对生成负载减少空泡与重复计算硬件后端统一抽象、渐进优化NVIDIA 与昇腾分别维护算子、布局与通信硬件利用率更高维护矩阵更大因此比较 DeepInfer、vLLM 和 SGLang 时不能只看 API 或模型数量。更合适的观察对象是目标模型、目标硬件和真实请求分布下的端到端 Pareto 前沿。四、昇腾适配不是简单兼容层DeepInfer 的昇腾路径在构建、运行时、拓扑、通信、算子和图执行六个层面都有独立实现。构建层通过DEEPINFER_ACCELERATORascend选择昇腾依赖与NpuExtension并加入asc_utils.cpp、graph_timer_asc.cpp等独立源码。运行时层使用torch_npu映射大量 CUDA 风格调用固定 TileLang 的ascend目标配置 HCCL 缓冲、超时和可扩展显存并固定主机内存。拓扑层通过/dev/davinci*、/sys/class/ubcore和npu-smi发现 NPU、UB 实体与 NUMA 关系。通信层分布式后端使用 HCCLDeepEP-Ascend 使用 HCCL / HCOMM、UBMEM、URMA 与 DeepJIT 完成 EP dispatch / combine。算子层Ascend TileLang 覆盖 KV、采样、DFlash、SWA、CP、视觉与结构化输出。FP8 scale 在昇腾上使用int16打包与 NVIDIA 的int32路径区分。图执行与测试层NPU Graph 支持捕获和 replay离线推理 CI 覆盖 2 节点 × 8 NPU。开源快照中各层已经形成清晰实现路径EPLB 的负载统计和重排计算已具备在线权重迁移可沿平台对称内存接口继续扩展。DeepSelect 已覆盖 BF16 TopKFP32 且需要返回 value 时由torch.topk保持功能兼容。NPU Graph Timer 已定义统一接口细粒度平台事件可按目标 CANN 工具链接入。FlashMLA 的昇腾实现聚焦 V4.1 稀疏 Attention 的 Prefill 与 Decode。完整 KV Cache 服务、P/D 编排和跨 DP 副本负载均衡属于外层服务系统开源引擎保留缓存初始化、保存钩子与 P/D 接入点。deepseek-recipe 的作用deepseek-recipe是 API 协议适配与 Prompt / Token 编解码组件负责 Chat Completions、Responses、Messages 转换、Prompt 渲染、Tokenizer 编解码、流式解析与工具调用。它不是推理引擎也不是 KV Cache、P/D 编排或副本负载均衡的补齐件。调用方仍需提供模型推理、工具执行与 HTTP 传输。因此DeepInfer 与 deepseek-recipe 是推理核心和表示层适配的关系而不是同一层能力的替代关系。五、昇腾 950 上的计算流与微架构优化DeepInfer 的昇腾实现不是把 CUDA kernel 逐行翻译成 NPU 代码而是围绕昇腾 AI Core 的异构执行资源重新安排数据流。重点资源包括CUBE矩阵计算。VECTOR向量计算、反量化、softmax 与结果整理。MTE2 / MTE3全局内存、L1 与片上缓冲之间的数据搬运。FixPipe矩阵累加结果的搬出与格式处理。UB统一向量缓冲承担向量计算与中间结果整理。L1 / L0C矩阵输入缓存与累加结果存储。优化覆盖 Ascend 950 的dav-3510共用核心架构。应用到 950PR 与 950DT 时算子能力和通信带宽会因具体型号、HDK、固件和组网配置而不同。Prefill计算受限目标是保持 CUBE 持续工作Prefill 一次处理较多 query Token矩阵计算量足够大。FlashMLA 昇腾实现为每个 AI Core 分配 64 个 query head采用 1 CUBE 2 VECTOR 的核内编组把 softmax、反量化、ND→NZ 和 FixPipe 搬运分摊到两个 VECTOR Core。流水使用“当前块 QK、上一块 PV、后续块预取”的交错方式使 CUBE 不必等待 softmax。skip-scale 只在 online softmax 的最大值偏移超过阈值时执行重缩放减少 FixPipe 和向量计算。Decode数据与搬运更敏感Decode 每步 query Token 少、KV 历史长更容易受 KV 读取、FP8 / FP4 反量化、索引 gather 和 FixPipe 影响。公开数据中Decode 为 360 TFLOPS利用率约 83%低于 Prefill 的 95%这一差异与 Decode 的结构性特征一致。因此实现将 KV 的 data 与 scale 按 Token 相邻存储减少 MTE2 请求使用 gather2 合并两个 Token 的搬运并在 UB 中完成 ND 到 NZ 的高吞吐转换。八项关键微架构优化优化原有问题实现方式作用ND → NZ 片上转换稀疏 KV 无法在 GM → L1 时顺路完成布局转换先进入 UB再由 SIMD VF 转换避免 L2 访存合并失效Token 交错 KV 布局data 与 scale 分离会增加搬运请求同一 Token 的 data / scale 相邻存放降低 Decode 的 MTE2 请求数和地址开销双 L1 bank 均衡单 bank 带宽可能限制矩阵供数将 Q、K、V 等张量分布到两个 L1 bank让 CUBE 输入供给更稳定昇腾专用 scale 布局GPU 与 NPU 的 scale ABI 不同FP8 scale 使用int16成对打包并按 MN-major 存放匹配昇腾硬件访问效率Skip-scale 与跨核同步每块都重缩放会增加 FixPipe 与向量运算SS buffer 与 cross-core flag 传递是否重缩放减少不必要的 UB 和 VECTOR 工作gather2 合并搬运大量细小 gather 请求会阻塞 MTE2一次请求搬运两个稀疏 Token降低 issue queue 压力指令缓存控制过度展开会让小容量 icache 频繁换入换出硬件循环、谨慎展开与函数内联保持 CUBE / VECTOR / SIMD VF 流水稳定FixPipe 感知流水输出 rescale 可能占用 L0C 并引入额外矩阵乘accumulator 放 UB按块执行 FixPipe 搬出与 VECTOR 累加减少 L0C 占用和额外计算这些优化的共同目标是减少 CUBE 等待供数、等待 softmax、等待搬运和等待同步的时间。矩阵乘峰值只代表理想条件下的计算能力真实效率取决于整条数据流能否持续供数。六、公开性能数据与来源公开性能数据来自 DeepInfer README 及其随仓库组件文档。下面分别列出计算吞吐和通信带宽不将组件实测直接等同于完整服务的端到端吞吐。FlashMLA 与 DeepGEMM-Ascend项目实测结果参考峰值利用率BF16 Dense GEMM431 TFLOPS432 TFLOPS99.8%FP8 × FP8 GEMM861 TFLOPS865 TFLOPS99.5%FP4 × FP4 GEMM1701 TFLOPS1730 TFLOPS98.3%稀疏 Attention Prefill410 TFLOPS432 TFLOPS约 95%稀疏 Attention Decode360 TFLOPS432 TFLOPS约 83%这里的 Prefill 数据应读作实测 410 TFLOPSBF16 参考峰值为 432 TFLOPS前者约为后者的 95%。410和432是两个独立的 TFLOPS 数值410 / 432不作为比值写法使用。DeepEP-Ascend 通信带宽配置dispatchcombineEP8373–375 GB/s345–347 GB/sEP64323–327 GB/s未列出EP128313–320 GB/s272–278 GB/sDeepEP-Ascend 的公开带宽数据对应 950DT、验证版 HDK 与外部 Clos 网络。不同 HDK、固件和网络拓扑下应重新测量不能直接外推为任意生产集群的带宽。数据来源DeepInfer READMEFlashMLA READMEFlashMLA 昇腾 Prefill 技术说明DeepGEMM-Ascend READMEDeepEP-Ascend READMEdeepseek-recipe README七、主要源码证据源码关键内容README.md说明 DeepInfer 支撑大规模生产服务开源版本聚焦离线推理与规模化服务核心能力deepinfer/engine/llm_engine.py缓存初始化、P/D 初始化、KV 保存与预填充钩子deepinfer/platform.pyHCCL、TileLangasc后端、torch_npu 映射、ACL 主机内存固定与 FP8 scale 类型deep_cluster/hardware/ascend.pydavinci 设备、ubcore sysfs 与 npu-smi 拓扑发现setup.py/build.shNpuExtension、昇腾专用 C 源码与构建选择deepinfer/model_executor/eplb/__init__.pyEPLB 负载统计、周期重排与平台 shuffle 扩展点deepinfer/npu_ops/deep_select.pyBF16 TopK 与 FP32 兼容路径csrc/graph_timer/graph_timer_asc.cppNPU Graph Timer 统一接口tests/models/test_offline_inference.py昇腾离线推理 CI 的 EP16 / CP8 配置deepinfer/model_executor/layers/tilelang_ops/asc/KV、采样、DFlash、SWA、CP、视觉与结构化输出算子八、核心结论DeepInfer 的技术价值不在于重新发明每一个并行名词而在于把模型结构、算子、通信、缓存语义和调度系统放进同一个优化闭环。对昇腾而言它已经形成从构建、运行时、拓扑、集合通信、TileLang 算子到 NPU Graph 的多层后端对 DeepSeek-V4.1-Flash 而言它利用固定模型带来的确定性减少长 Prompt、MoE 热点、逐 Token Decode、缓存重放和数据搬运形成的等待。公开数据表明FlashMLA、DeepGEMM-Ascend 和 DeepEP-Ascend 在目标配置下已经达到较高组件效率。下一步评价应聚焦目标模型、目标 HDK 和真实请求分布下的完整服务链路包括首 Token 延迟、每 Token 输出时间、吞吐、缓存命中与服务稳定性。术语速查DSADeepSeek Sparse Attention按索引只关注部分 KV Token。EPExpert Parallel将 MoE 专家分布到不同 rank。EPLBExpert Parallel Load Balancing根据负载重排专家。CPContext Parallel将长上下文切分到不同 rank。P/D 分离Prefill 与 Decode 分别部署和调度。HCCL / HCOMM昇腾集合通信与底层通信组件。URMA / UBMEMDeepEP-Ascend 使用的远端内存访问与统一总线内存能力。TFLOPS每秒万亿次浮点运算。
完整交互版https://besproma.github.io/deepinfer-ascend-analysis/代码快照774ce24分析日期2026-10-21分析范围并行调度、模型级优化、昇腾适配、芯片内数据流、源码证据与公开性能数据DeepInfer 是 DeepSeek 面向 DeepSeek-V4.1-Flash 深度定制的高性能推理框架。它与通用推理框架的主要差别不是又增加了一套模型接入接口而是把模型结构、稀疏 Attention、MoE 路由、KV 规则和服务调度放在同一个设计闭环中并分别实现 NVIDIA 与昇腾后端。本文重点回答三个问题DeepInfer 为什么要做这些优化昇腾适配到底落在哪些工程层以及公开性能数据能够说明什么。一、为什么要重新组织推理路径一次大模型请求看似只是“输入 Prompt、输出 Token”实际会经过一条很长的关键路径Prefill 阶段处理整段 Prompt并产生首 Token。MoE 层根据路由把 Token 分发到不同专家再汇合结果。Decode 阶段逐个 Token 自回归生成。历史上下文在 KV Cache 中复用缓存未命中时还要恢复。CPU 侧持续完成请求检查、批处理拼装和元数据准备。这条链路的耗时往往不只来自矩阵计算。长 Prompt 可能只落在一组设备上MoE 的热点专家可能让最忙 rank 决定整轮时间Decode 每一步都很小却要重复读取权重和启动 kernelCPU 排批、集合通信和中间张量写回还会制造额外等待。DeepInfer 的核心思路是让模型已经知道的形状与规则进入执行系统提前消除这些等待。关键路径常见问题DeepInfer 对应机制目标长上下文 Prefill单条请求难以拆到更多设备Balanced CP缩短首 Token 关键路径MoE 专家路由热点专家与跨 rank 分发形成等待EP、DeepEP、EPLB扩大专家池并压低通信与尾部负载逐 Token Decode串行步数多小 batch 算力利用率低DSpark一次验证多个候选 TokenKV 缓存恢复已失效历史仍被完整重放SWA bounded replay只处理当前输出需要的最短后缀CPU / NPU 协同设备执行期间 CPU 在准备下一批异步批调度让排批与设备执行重叠MoE 数据路径多个 kernel 反复写回和读取中间结果MegaMoE把分发、GEMM、SwiGLU、汇合接成连续数据流这些机制并非每一项都由 DeepInfer 首次提出。它的差异化在于围绕一个固定模型把并行方式、通信布局、缓存语义与硬件数据流一起调整。二、七项关键技术如何减少等待1. Balanced CP让一条长请求也能多卡并行传统的数据并行主要增加同时处理的请求数量却不会自动拆开单条长 Prompt。当一条请求的上下文很长时它仍然可能由一组设备独自承担某个 rank 分到更多 Token其他 rank 就要等待形成木桶效应。Balanced CP 按上下文切分一条请求并尽量平衡各 rank 的 Token 数和计算量。它扩大单请求可用的计算资源同时缩小不同 rank 的完成时间差因此优化的是首 Token 阶段的尾延迟。2. MoE EP 与 DeepEP把专家容量扩到集群MoE 模型包含大量专家。专家并行把不同专家分布到不同 rank从而扩大可容纳的权重和可用算力。但 Token 必须先通过 dispatch 送到目标专家计算后再通过 combine 返回原 rank通信路径处理不好新增算力会被跨卡等待抵消。DeepEP-Ascend 负责昇腾上的分发与汇合并接入 HCCL、HCOMM、UBMEM 和 URMA 等通信能力。它的价值来自两件事同时成立专家计算扩展到更多设备跨设备通信成本得到控制。公开验证覆盖 EP128。这里的 EP128 表示 128 个专家并行 rank不是 128 个专家。公开数据中dispatch 峰值约为 313–320 GB/s。3. EPLB让专家布局跟随真实热点专家铺到更多设备之后总量通常不再是主要矛盾局部热点才是。某些专家承接的 Token 明显更多对应 rank 就会成为新的长尾。EPLB 先观察真实 Token 的负载再计算新的专家布局并通过重排压缩最忙 rank 的排队时间。当前实现已经覆盖负载统计、方案求解和周期重排在线权重迁移保留平台对称内存接口便于继续利用昇腾的统一总线能力。4. DSpark把一步一 Token 改成一次核验多个标准自回归 Decode 每轮只前进一个 Token。生成序列较长时串行轮次很多小 batch 下设备算力难以铺满但每轮仍要承担权重读取、kernel 启动和调度成本。DSpark 先生成多个轻量候选 Token再由主模型一次核验。主模型接受正确前缀后一次迭代可以前进多个 Token。实际收益取决于候选接受率、核验成本和额外调度开销。昇腾算子目录中已经包含 DFlash、采样和相关调度路径说明该机制已进入硬件后端实现而不只停留在模型语义层。5. SWA bounded replay不要重算已经失效的历史滑动窗口注意力只关注最近的上下文窗口。缓存恢复时如果仍按通用路径完整重放历史窗口之外的 Token 虽然已经不影响当前输出却还会消耗 KV 读取、数据搬运和 Prefill 计算。SWA bounded replay 只重放覆盖当前注意窗口和最大回溯长度的必要后缀。V4.1 默认保留 140 个 Token具体收益取决于窗口配置、缓存复用率和命中率。6. 异步批调度下一批在上一批执行时准备好如果 CPU 和 NPU 严格轮流工作每一轮都要先完成请求检查、batch 拼装和元数据准备再交给设备执行。Decode 单批执行时间较短时这部分 CPU 等待会占据明显比例。异步调度让 CPU 在 NPU 执行当前 batch 时准备下一批。理想情况下一轮耗时更接近“CPU 准备时间”和“NPU 执行时间”中较长的一项而不是两者相加。NPU Graph 还能通过捕获与 replay 降低 kernel 启动开销。7. MegaMoE把多次 kernel 往返接成一条数据流dispatch、grouped GEMM、SwiGLU 和 combine 如果分别作为独立 kernel 执行每一步都要把中间结果写回再由下一步读取。这会增加访存、缓冲区占用、kernel 启动和同步点。MegaMoE 让通信布局、专家矩阵形状和中间结果传递彼此匹配减少中间张量的往返并创造通信与计算重叠的机会。它优化的是整条 MoE 数据路径而不是单个 GEMM 的峰值。三、DeepInfer 与 vLLM、SGLang 的关系DeepInfer 最初主要基于 vLLM 构建同时吸收了 SGLang、TensorRT-LLM 等框架的设计思路。它并非从零重写所有通用能力也不是替代所有推理框架。与 vLLM、SGLang 的主要差异可以概括为通用框架优先覆盖更多模型、硬件和接口DeepInfer 则只围绕 DeepSeek-V4.1-Flash 做单模型垂直优化。固定模型带来了更强的协同设计空间算子可以围绕固定 head 形状、稀疏模式和精度格式做极致调优。通信可以理解 MoE 路由和 GEMM 时序而不只提供通用 all-to-all。调度可以把 Draft Token、SWA 重放、缓存布局和 NPU Graph 一起考虑。昇腾后端可以针对 CUBE、VECTOR、L1、UB 和 FixPipe 重新组织数据流。代价同样明确模型泛化能力较弱版本耦合更强每代硬件都需要维护独立实现。设计维度通用推理框架常见策略DeepInfer 策略技术含义模型范围多模型、多模态、广泛兼容当前聚焦 DeepSeek-V4.1-Flash可以做模型级协同跨模型复用需重新适配Attention通用 PagedAttention / FlashAttentionDSA 稀疏索引、FP8 / FP4 KV、固定 head 形状稀疏场景优化空间更大非目标模型不适用MoE 通信通用 all-to-all 与集合通信DeepEP 与 MegaMoE面向专家 GEMM 的通信和计算协同执行调度通用连续批处理异步调度、DSpark、SWA 有界重放针对生成负载减少空泡与重复计算硬件后端统一抽象、渐进优化NVIDIA 与昇腾分别维护算子、布局与通信硬件利用率更高维护矩阵更大因此比较 DeepInfer、vLLM 和 SGLang 时不能只看 API 或模型数量。更合适的观察对象是目标模型、目标硬件和真实请求分布下的端到端 Pareto 前沿。四、昇腾适配不是简单兼容层DeepInfer 的昇腾路径在构建、运行时、拓扑、通信、算子和图执行六个层面都有独立实现。构建层通过DEEPINFER_ACCELERATORascend选择昇腾依赖与NpuExtension并加入asc_utils.cpp、graph_timer_asc.cpp等独立源码。运行时层使用torch_npu映射大量 CUDA 风格调用固定 TileLang 的ascend目标配置 HCCL 缓冲、超时和可扩展显存并固定主机内存。拓扑层通过/dev/davinci*、/sys/class/ubcore和npu-smi发现 NPU、UB 实体与 NUMA 关系。通信层分布式后端使用 HCCLDeepEP-Ascend 使用 HCCL / HCOMM、UBMEM、URMA 与 DeepJIT 完成 EP dispatch / combine。算子层Ascend TileLang 覆盖 KV、采样、DFlash、SWA、CP、视觉与结构化输出。FP8 scale 在昇腾上使用int16打包与 NVIDIA 的int32路径区分。图执行与测试层NPU Graph 支持捕获和 replay离线推理 CI 覆盖 2 节点 × 8 NPU。开源快照中各层已经形成清晰实现路径EPLB 的负载统计和重排计算已具备在线权重迁移可沿平台对称内存接口继续扩展。DeepSelect 已覆盖 BF16 TopKFP32 且需要返回 value 时由torch.topk保持功能兼容。NPU Graph Timer 已定义统一接口细粒度平台事件可按目标 CANN 工具链接入。FlashMLA 的昇腾实现聚焦 V4.1 稀疏 Attention 的 Prefill 与 Decode。完整 KV Cache 服务、P/D 编排和跨 DP 副本负载均衡属于外层服务系统开源引擎保留缓存初始化、保存钩子与 P/D 接入点。deepseek-recipe 的作用deepseek-recipe是 API 协议适配与 Prompt / Token 编解码组件负责 Chat Completions、Responses、Messages 转换、Prompt 渲染、Tokenizer 编解码、流式解析与工具调用。它不是推理引擎也不是 KV Cache、P/D 编排或副本负载均衡的补齐件。调用方仍需提供模型推理、工具执行与 HTTP 传输。因此DeepInfer 与 deepseek-recipe 是推理核心和表示层适配的关系而不是同一层能力的替代关系。五、昇腾 950 上的计算流与微架构优化DeepInfer 的昇腾实现不是把 CUDA kernel 逐行翻译成 NPU 代码而是围绕昇腾 AI Core 的异构执行资源重新安排数据流。重点资源包括CUBE矩阵计算。VECTOR向量计算、反量化、softmax 与结果整理。MTE2 / MTE3全局内存、L1 与片上缓冲之间的数据搬运。FixPipe矩阵累加结果的搬出与格式处理。UB统一向量缓冲承担向量计算与中间结果整理。L1 / L0C矩阵输入缓存与累加结果存储。优化覆盖 Ascend 950 的dav-3510共用核心架构。应用到 950PR 与 950DT 时算子能力和通信带宽会因具体型号、HDK、固件和组网配置而不同。Prefill计算受限目标是保持 CUBE 持续工作Prefill 一次处理较多 query Token矩阵计算量足够大。FlashMLA 昇腾实现为每个 AI Core 分配 64 个 query head采用 1 CUBE 2 VECTOR 的核内编组把 softmax、反量化、ND→NZ 和 FixPipe 搬运分摊到两个 VECTOR Core。流水使用“当前块 QK、上一块 PV、后续块预取”的交错方式使 CUBE 不必等待 softmax。skip-scale 只在 online softmax 的最大值偏移超过阈值时执行重缩放减少 FixPipe 和向量计算。Decode数据与搬运更敏感Decode 每步 query Token 少、KV 历史长更容易受 KV 读取、FP8 / FP4 反量化、索引 gather 和 FixPipe 影响。公开数据中Decode 为 360 TFLOPS利用率约 83%低于 Prefill 的 95%这一差异与 Decode 的结构性特征一致。因此实现将 KV 的 data 与 scale 按 Token 相邻存储减少 MTE2 请求使用 gather2 合并两个 Token 的搬运并在 UB 中完成 ND 到 NZ 的高吞吐转换。八项关键微架构优化优化原有问题实现方式作用ND → NZ 片上转换稀疏 KV 无法在 GM → L1 时顺路完成布局转换先进入 UB再由 SIMD VF 转换避免 L2 访存合并失效Token 交错 KV 布局data 与 scale 分离会增加搬运请求同一 Token 的 data / scale 相邻存放降低 Decode 的 MTE2 请求数和地址开销双 L1 bank 均衡单 bank 带宽可能限制矩阵供数将 Q、K、V 等张量分布到两个 L1 bank让 CUBE 输入供给更稳定昇腾专用 scale 布局GPU 与 NPU 的 scale ABI 不同FP8 scale 使用int16成对打包并按 MN-major 存放匹配昇腾硬件访问效率Skip-scale 与跨核同步每块都重缩放会增加 FixPipe 与向量运算SS buffer 与 cross-core flag 传递是否重缩放减少不必要的 UB 和 VECTOR 工作gather2 合并搬运大量细小 gather 请求会阻塞 MTE2一次请求搬运两个稀疏 Token降低 issue queue 压力指令缓存控制过度展开会让小容量 icache 频繁换入换出硬件循环、谨慎展开与函数内联保持 CUBE / VECTOR / SIMD VF 流水稳定FixPipe 感知流水输出 rescale 可能占用 L0C 并引入额外矩阵乘accumulator 放 UB按块执行 FixPipe 搬出与 VECTOR 累加减少 L0C 占用和额外计算这些优化的共同目标是减少 CUBE 等待供数、等待 softmax、等待搬运和等待同步的时间。矩阵乘峰值只代表理想条件下的计算能力真实效率取决于整条数据流能否持续供数。六、公开性能数据与来源公开性能数据来自 DeepInfer README 及其随仓库组件文档。下面分别列出计算吞吐和通信带宽不将组件实测直接等同于完整服务的端到端吞吐。FlashMLA 与 DeepGEMM-Ascend项目实测结果参考峰值利用率BF16 Dense GEMM431 TFLOPS432 TFLOPS99.8%FP8 × FP8 GEMM861 TFLOPS865 TFLOPS99.5%FP4 × FP4 GEMM1701 TFLOPS1730 TFLOPS98.3%稀疏 Attention Prefill410 TFLOPS432 TFLOPS约 95%稀疏 Attention Decode360 TFLOPS432 TFLOPS约 83%这里的 Prefill 数据应读作实测 410 TFLOPSBF16 参考峰值为 432 TFLOPS前者约为后者的 95%。410和432是两个独立的 TFLOPS 数值410 / 432不作为比值写法使用。DeepEP-Ascend 通信带宽配置dispatchcombineEP8373–375 GB/s345–347 GB/sEP64323–327 GB/s未列出EP128313–320 GB/s272–278 GB/sDeepEP-Ascend 的公开带宽数据对应 950DT、验证版 HDK 与外部 Clos 网络。不同 HDK、固件和网络拓扑下应重新测量不能直接外推为任意生产集群的带宽。数据来源DeepInfer READMEFlashMLA READMEFlashMLA 昇腾 Prefill 技术说明DeepGEMM-Ascend READMEDeepEP-Ascend READMEdeepseek-recipe README七、主要源码证据源码关键内容README.md说明 DeepInfer 支撑大规模生产服务开源版本聚焦离线推理与规模化服务核心能力deepinfer/engine/llm_engine.py缓存初始化、P/D 初始化、KV 保存与预填充钩子deepinfer/platform.pyHCCL、TileLangasc后端、torch_npu 映射、ACL 主机内存固定与 FP8 scale 类型deep_cluster/hardware/ascend.pydavinci 设备、ubcore sysfs 与 npu-smi 拓扑发现setup.py/build.shNpuExtension、昇腾专用 C 源码与构建选择deepinfer/model_executor/eplb/__init__.pyEPLB 负载统计、周期重排与平台 shuffle 扩展点deepinfer/npu_ops/deep_select.pyBF16 TopK 与 FP32 兼容路径csrc/graph_timer/graph_timer_asc.cppNPU Graph Timer 统一接口tests/models/test_offline_inference.py昇腾离线推理 CI 的 EP16 / CP8 配置deepinfer/model_executor/layers/tilelang_ops/asc/KV、采样、DFlash、SWA、CP、视觉与结构化输出算子八、核心结论DeepInfer 的技术价值不在于重新发明每一个并行名词而在于把模型结构、算子、通信、缓存语义和调度系统放进同一个优化闭环。对昇腾而言它已经形成从构建、运行时、拓扑、集合通信、TileLang 算子到 NPU Graph 的多层后端对 DeepSeek-V4.1-Flash 而言它利用固定模型带来的确定性减少长 Prompt、MoE 热点、逐 Token Decode、缓存重放和数据搬运形成的等待。公开数据表明FlashMLA、DeepGEMM-Ascend 和 DeepEP-Ascend 在目标配置下已经达到较高组件效率。下一步评价应聚焦目标模型、目标 HDK 和真实请求分布下的完整服务链路包括首 Token 延迟、每 Token 输出时间、吞吐、缓存命中与服务稳定性。术语速查DSADeepSeek Sparse Attention按索引只关注部分 KV Token。EPExpert Parallel将 MoE 专家分布到不同 rank。EPLBExpert Parallel Load Balancing根据负载重排专家。CPContext Parallel将长上下文切分到不同 rank。P/D 分离Prefill 与 Decode 分别部署和调度。HCCL / HCOMM昇腾集合通信与底层通信组件。URMA / UBMEMDeepEP-Ascend 使用的远端内存访问与统一总线内存能力。TFLOPS每秒万亿次浮点运算。
平方复杂度无法被常规优化消除｜TaoToken 统一 Key 通道实测)
