作者 | 林绮蓓
编辑 | Tina
9 月 30 日,DeepSeek 宣布开源面向华为昇腾算力平台的一组基础设施组件,涵盖 TileLang 高级语言编译工具,以及 DeepGEMM、DeepEP、TileKernels、FlashMLA 和 DeepSelect 的昇腾平台实现。据官方公告,这些组件与此前面向英伟达平台开源的组件一一对应,覆盖矩阵运算、跨设备通信、常规向量计算与访存、稀疏注意力和数据筛选等功能。
DeepSeek 同时披露,团队与华为共同推进了基于昇腾 950 的 128 卡超节点方案,并对计算和通信进行了联合优化。对开发者而言,此次发布提供了可用于昇腾环境的底层实现、接口和测试代码。其中,部分组件沿用英伟达版本上层接口,便于已有调用代码迁移;具体硬件要求、数据格式与功能支持范围,则需要按组件分别确认。
TileLang 昇腾后端提供哪些能力
在这组项目中,TileLang 负责提供算子开发所用的语言和编译能力,其余五类组件提供可调用的计算与通信实现。开发者既可以使用现成组件,也可以通过 TileLang 编写或修改内核,以适配自己的模型与计算任务。
TileLang 是用于编写高性能计算内核的领域专用语言,采用接近 Python 的语法,其编译基础设施建立在 TVM 之上。项目已于 9 月 30 日在主仓库中公布对昇腾 950 NPU 的支持,包括原生代码生成、自动调度和同步等能力。
据 DeepSeek 介绍,TileLang 已承载 DeepSeek V4 系列模型训练中大部分算子的实现;目前,其训练中使用的每一个 TileLang 算子,在昇腾上均有对应的高性能实现。官方称,昇腾版本通过封装 Ascend C 底层指令,提供高级语言编程方式,同时保留针对硬件进行优化的能力。
具体而言,昇腾后端复用了 TileLang 的共享前端,并针对 Cube 矩阵计算单元和 Vector 向量计算单元,实现专用的编译转换、调度、同步和代码生成。开发者仍需描述数据类型、分块大小和计算操作,编译器则可承担部分数据搬运、流水线安排及单元间同步工作。这为自定义算子开发提供了更高层的表达方式。不过,昇腾后端也包含专用语言接口和存储操作,跨平台适配仍需结合目标硬件的执行方式进行调整。
五类计算与通信组件分别解决什么问题
此次开源的五类计算与通信组件,分别覆盖矩阵运算、跨设备通信、稀疏注意力、通用算子和数据筛选,为模型训练与推理提供可复用的基础能力。
DeepGEMM-Ascend 用于加速模型中的矩阵乘法,支持 BF16、FP8、FP4 等精度。项目声明其 API 与 DeepGEMM 兼容,开发者可沿用既有接口和开发流程;不过,昇腾版本的量化缩放因子存储格式与英伟达版本存在差异,迁移时仍需适配数据布局。
DeepEP Ascend 负责跨设备通信,提供混合专家模型(MoE)专家并行所需的 all-to-all 操作,将数据分发至不同设备上的专家,并汇总计算结果。其公开 buffer API 与英伟达版本对齐,支持 FP8 数据分发;流水线并行、上下文并行和数据并行相关的部分通信能力仍在开发中。
FlashMLA 用于加速稀疏注意力计算,此次昇腾支持覆盖预填充和解码,分别对应处理输入上下文和后续逐步生成 token 的阶段。目前,仓库中仍有部分融合内核及稠密注意力内核仅支持 CUDA。
TileKernels 提供基于 TileLang 编写的常用算子,覆盖 MoE 路由、量化、位置编码等训练和推理操作。新增的昇腾后端支持运行时自动选择,使开发者能够通过同一组 Python API 调用英伟达 GPU 和华为 NPU 上的相应实现。
DeepSelect 提供 TopK 数据筛选算子,用于从候选数据中选出得分最高的 K 项,面向 DeepSeek 稀疏注意力及采样器相关场景。其昇腾实现目前仅支持 BF16 输入,CUDA 实现还支持 FP32 输入,开发者需根据具体场景核对平台支持范围。
对于已有相关组件调用代码的团队,API 对齐和多后端支持提供了复用既有代码的条件;对于需要自定义算子的团队,TileLang 昇腾后端及配套实现则提供了开发入口和优化参考。
此次公开文档列出的目标硬件和验证环境主要围绕昇腾 950 展开:TileKernels 的昇腾后端要求 CANN 9.2.0 及以上版本;DeepEP Ascend 则说明,现有测试尚未验证其内核对其他昇腾代际或 CANN 版本的支持情况。实际接入时仍需核对目标硬件、软件版本、数据类型和算子覆盖范围,并在自身负载上验证正确性与性能。此次发布提供了从算子编写到计算、通信执行的多层组件;完整模型的部署效果,还取决于这些组件与上层框架、模型实现及集群环境的集成情况。
同样是计算、通信和注意力,此次有何不同?
此前,华为与硅基流动的一篇论文《Serving Large Language Models on Huawei CloudMatrix384》 ,介绍了在昇腾 910 系列超节点上运行 DeepSeek-R1 的方案。论文既讨论如何分配预填充、解码和缓存所需的资源,也涉及矩阵计算、专家通信和注意力算子的优化。此次 DeepSeek 公布的内容则主要面向昇腾 950,将编译工具和计算、通信基础库开放出来,供开发者调用、修改并接入自己的模型与框架。
在具体计算上,论文中重点讨论了 INT8 量化在矩阵运算和数据传输中的应用;而此次 DeepGEMM-Ascend 提供 BF16、FP8、FP4 等矩阵计算实现,DeepEP-Ascend 支持 FP8 数据分发。两者都涉及通过低精度计算或传输减少开销,但支持的数据格式有所不同。
对于专家并行中的数据分发与结果汇总,论文介绍了 FusedDispatch、FusedCombine 等融合算子,并利用微批次流水线安排计算与通信;此次除了开放异步通信接口,还在 MegaMoE 中将专家分发、矩阵计算、激活和结果汇总融合执行。
注意力部分的变化,与模型本身的变化紧密相关。针对 R1 的 MLA,论文将归一化、线性变换、位置编码等步骤合并到 MLAProlog 等算子中,并调整 KV 缓存布局,以减少零散操作和格式转换带来的开销。V4 则沿序列维度压缩 KV 缓存,并在部分注意力计算中筛选需要访问的内容,因而增加了索引、筛选和缓存管理方面的适配需求。此次 FlashMLA 提供稀疏注意力的预填充与解码内核,DeepSelect 提供 TopK 筛选能力,分别支持其中的计算环节。
算子开发方式也有所变化。论文介绍了 CANN 软件栈及 Ascend C 自定义算子开发能力;此次 TileLang 昇腾后端提供较高层的编程接口,由编译器完成相应的代码生成、调度和同步。TileKernels 则将 MoE 路由、量化、位置编码等常用操作封装为算子库,便于模型代码调用。
这些工具仍需与 CANN 等底层软件配合使用。它们与论文中的部分模块解决相近的问题,但公开资料尚未明确两者的代码继承关系,也尚不能将此次发布认定为 CloudMatrix-Infer 的后续版本。
从模型适配到联合优化:DeepSeek 与华为的技术交集
沿公开信息回看,DeepSeek 模型与昇腾平台的交集,经历了云端推理服务适配、超节点系统优化、算子开发工具适配等进展。参与其中的既有 DeepSeek 和华为,也有硅基流动等服务提供方。
2025 年 1 月,硅基流动与华为云上线基于昇腾的 DeepSeek-V3、R1 推理服务。硅基流动在公告中明确表示,模型适配得到了 DeepSeek 与华为云的支持。对应用开发者而言,当时提供的主要入口是云端 API:无需自行部署模型,即可调用运行在昇腾环境中的推理服务。这是三方在模型、算力与服务交付上的公开交集。
同年 6 月,华为与硅基流动通过 CloudMatrix384 论文,进一步披露运行 DeepSeek-R1 的超节点推理方案,将相关实践从服务上线延伸到硬件互连、资源组织和软件优化的技术说明。
2025 年 9 月,伴随 DeepSeek-V3.2-Exp 发布,昇腾完成 TileLang 适配,并开源稀疏注意力及索引算子的相关实现。2026 年 4 月,华为云在 V4 预览版发布当天宣布完成适配,通过注意力、缓存管理和融合算子等优化,支持百万 token 上下文推理。
此次,DeepSeek 明确披露了与华为共同推进昇腾 950 的 128 卡超节点方案,以及计算和通信方面的联合优化。从早期的云端推理服务,到超节点系统和算子开发工具,相关工作逐步延伸到模型执行的不同层次。此次基础设施组件的开放,也为开发者在昇腾平台开展模型适配和性能优化提供了可复用的代码。
附开源项目链接:
TileLang:
https://github.com/tile-ai/tilelang
DeepGEMM Ascend:
https://github.com/deepseek-ai/DeepGEMM-Ascend
DeepEP Ascend:
https://github.com/deepseek-ai/DeepEP-Ascend
TileKernels:
https://github.com/deepseek-ai/TileKernels
FlashMLA:
https://github.com/deepseek-ai/FlashMLA
DeepSelect: