2026 年 10 月 01 日——9 月 30 日,DeepSeek 宣布正式开源面向华为昇腾算力平台的基础设施组件,涵盖 TileLang 编译工具、四个计算库与分布式通信库 DeepEP 的昇腾版本,与此前面向英伟达平台的开源组件一一对应;官方文档显示,DeepGEMM-Ascend 的 Dense GEMM 在昇腾 950DT 上利用率最高达到硬件极限的 99.8%,DeepEP-Ascend 的 Dispatch 带宽最高达 375 GB/s。
这套组件的特别之处在于「同构」而非「移植」:TileKernels 等库提供同一套 Python API,运行时自动选择 NVIDIA 或昇腾后端;商用固件包预计 10 月 15 日前后由华为发布,届时公开用户可在正式环境验证官方给出的性能数据。
2026-10-01 · DeepSeek · 华为昇腾 · 开源
一、概览
据凤凰网科技等媒体报道,DeepSeek 于 9 月 30 日宣布开源面向华为昇腾算力平台的基础设施组件,涵盖 TileLang 高级语言编译工具、计算库及分布式通信库,所有组件与此前面向英伟达平台的开源组件一一对应。DeepSeek 官方 GitHub 仓库同日上线了 DeepGEMM-Ascend、DeepEP-Ascend、TileKernels、DeepSelect、DeepJIT 等仓库,均已开放源代码。
这并非 DeepSeek 第一次将训练基础设施开源。2025 年 2 月的开源周上,DeepSeek 曾连续开源 FlashMLA、DeepEP、DeepGEMM 等英伟达平台组件,覆盖注意力内核、MoE 专家并行通信、矩阵乘法等大模型训练的关键路径。本次发布相当于把同一条软件栈整体搬到昇腾平台,并以一一对应的方式开放。
华为方面的配合同步公开:据官方公告介绍,华为在 CANN 社区开源了配套的部署与训练实践,附 DeepSeek V4.1-Flash 每卡输出吞吐实测数据;媒体普遍注意到 DeepSeek 在公告中感谢华为「毫无保留的大力支持」。
二、核心信息
2.1 六大组件与英伟达平台一一对应
本次开源的组件按「编译工具、计算库、通信库」三层组织,与英伟达平台组件形成明确映射:
| 昇腾平台组件 | 英伟达平台对应组件 | 类型 | 功能定位 |
|---|---|---|---|
| TileLang(昇腾版) | TileLang(NVIDIA 版) | 编译工具 | 封装昇腾 Ascend C 底层指令,提供高级语言算子开发方式 |
| DeepGEMM-Ascend | DeepGEMM | 计算库 | BF16/FP8/FP4 矩阵乘法内核库 |
| TileKernels | TileKernels(同库双后端) | 计算库 | MoE 路由、Engram、量化、mHC 等数十个算子 |
| FlashMLA | FlashMLA | 计算库 | MLA 注意力高效内核 |
| DeepSelect | DeepSelect | 计算库 | DeepSeek 稀疏注意力(DSA)TopK 内核与采样器 |
| DeepEP-Ascend | DeepEP | 通信库 | MoE 专家并行 all-to-all 通信,支持 FP8 dispatch |
| DeepJIT | 无独立对应 | 配套设施 | xPU 内核运行时 JIT 编译库,为上述组件提供编译底座 |
从映射关系可以看出两个设计取向。第一,昇腾版本不是另起炉灶的新 API,而是尽量复用英伟达版本的接口形态:DeepGEMM-Ascend 与 DeepGEMM 完全 API 兼容且共用同一个 Python 包名 deep_gemm,DeepEP-Ascend 的公开 buffer API 与 NVIDIA 版 DeepEP 对齐,用户沿用原有代码与开发流程即可切换硬件。第二,配套的 DeepJIT 把内核编译挪到运行时按需进行,配合 JIT 缓存机制,使同一份代码可以在不同型号的 xPU 上完成本地编译。
2.2 TileLang 与 TileKernels:一套代码,双后端运行
快科技报道称,这套技术方案已在英伟达平台完成落地验证,DeepSeek V4 系列模型训练的绝大多数算子均基于 TileLang 实现。TileLang 是一个支持多硬件后端的领域专用语言(DSL),基于 Pythonic 语法与 TVM 编译器基础设施构建;其昇腾版本对昇腾 Ascend C 底层指令进行封装,官方称在提供高级语言编程方式的同时不损失硬件性能。GitHub 上 tile-ai 组织的 tilelang-ascend 仓库自 2025 年 9 月底开源以来持续迭代,2026 年 4 月已发布 DeepSeek V4 内核示例,编译后端支持 Ascend C & PTO 与 AscendNPU IR 两条技术路线。
TileKernels 则是这趟「双后端」路线的直接证明。这个包含数十个深度优化算子的库在 9 月 30 日的更新中加入了华为昇腾支持:沿袭 NVIDIA 路径,内核现在提供第二个后端,在运行时自动选择,同一套 Python API 可以同时运行在 NVIDIA GPU 和华为 NPU 上。其算子覆盖混合专家路由(MoE Routing)、Engram 门控(融合 RMSNorm,含前向、反向与权重梯度归约)、逐 token/逐块/逐通道的 FP8/FP4 量化(含融合 SwiGLU 与量化)、流形超连接 mHC(含 Sinkhorn 归一化)以及 RoPE 等。官方说明中称,大多数算子的性能接近硬件的计算吞吐或内存带宽上限,全部算子已用于 DeepSeek 内部训练与推理任务。
对开发者而言,「一套代码、双后端」意味着算子层的硬件差异被 DSL 和运行时吸收:写算子的人不再需要分别维护 CUDA 与 CANN 两套实现,硬件切换在部署侧完成。这也是本次开源与以往「单独移植某个算子库」的做法在工程理念上的主要差别。
2.3 DeepGEMM-Ascend:GEMM 利用率最高 99.8%
DeepGEMM-Ascend 是 DeepGEMM 在昇腾平台上的实现,9 月 30 日初始发布即支持 Ascend 950 系列设备。它对昇腾矩阵乘加原语(MAD)做了轻量抽象,隐藏分形矩阵布局、对齐约束、地址计算等底层细节,并采用稀疏数据加载、基于协程的流水线等昇腾特有优化。除标准 GEMM 外,还支持 MQA logits(DeepSeek Lightning Indexer 的打分内核)与 MegaMoE——后者将 EP dispatch、两组分组 GEMM、SwiGLU 与 combine 融合为单个内核。
官方性能数据在昇腾 950DT(CANN 9.20)上以 bench_msprof、冷 L2 缓存测得,测试形状沿用 DeepGEMM 测试集、覆盖 DeepSeek 系列模型的训练与推理负载。Dense GEMM 各精度的利用率如下:
| 精度 | 矩阵形状(M×N×K) | 实测算力(TFLOPS) | 硬件极限(TFLOPS) | 利用率 |
|---|---|---|---|---|
| FP4×FP4 | 4096×7168×16384 | 1701 | 1730 | 98.3% |
| FP8×FP4 | 4096×7168×16384 | 861 | 865 | 99.5% |
| FP8×FP8 | 4096×7168×16384 | 861 | 865 | 99.5% |
| BF16×BF16 | 4096×7168×16384 | 431 | 432 | 99.8% |
官方文档称,DeepGEMM Ascend 可以在多种矩阵形状上达到硬件极限性能;MQA logits 内核为 FIX 管线受限而非计算受限,FIX 管线利用率达 99%;mHC prenorm GEMM 在较大 M 规模下内存带宽达 3463 GB/s,接近打满 HBM 写带宽。这些数字共同支撑了「组件性能已接近硬件上限」的官方表态。
值得注意的一个实现细节是缩放因子格式差异:昇腾平台沿 K 维每对 UE8M0 缩放因子打包为一个 int16 并按 MN 主序存储,与 NVIDIA 平台的格式不同。这提示两套硬件在 FP8/FP4 低精度路径上的底层约定并不一致,「API 兼容」吸收的是接口层差异,数据布局层的差异仍由各 Ascend 版本库在内部处理。
2.4 DeepEP-Ascend:通信带宽逼近硬件上限
DeepEP-Ascend 面向昇腾 NPU 提供机器学习训练与推理的高性能通信能力,核心是 MoE dispatch/combine 的专家并行(EP)all-to-all 操作,支持 FP8 dispatch 与延迟 epilogue,另提供流水线并行(PP)、面向上下文与数据并行的 Bucket 集合通信以及 Engram 远端内存访问等原语(后三项标注为开发中)。内核层面使用 HCCL/HCOMM、UBMEM 和 URMA 完成通信,并经 DeepJIT 在运行时编译。
官方实测同样基于昇腾 950DT 与 CANN 9.2.0,配置为每 rank 16384 tokens、hidden size 7168、top-6 路由、256 专家,Dispatch 采用 FP8、Combine 采用 BF16:
| EP 规模 | Dispatch(GB/s) | Combine(GB/s) |
|---|---|---|
| EP8 | 373–375 | 345–347 |
| EP16 | 348–352 | 338–341 |
| EP32 | 335–340 | 320–324 |
| EP64 | 323–327 | 294–298 |
| EP128 | 313–320 | 272–278 |
官方文档称,EP32 及以下规模时 Dispatch 达到物理载荷带宽上限的约 90% 至 95%;更大的 EP 规模与 Combine 仍在优化中,Combine 额外承担本地归约开销、并与 URMA 存在 HBM 争用,华为后续固件改进计划缓解这一争用。
2.5 华为配套与商用进度
两个背景信息界定了这批数据的适用范围。其一,上述性能均基于华为向 DeepSeek 提供的 PoC HDK(手工配置)测得,并非公开发布的固件版本;包含满带宽运行所需配置的华为 Q3 商用 HDK(Atlas 850E)计划于 2026 年 10 月中旬(约 10 月 15 日)通过华为支持网站公开,具体日期以华为发布节奏为准。其二,据媒体报道,官方并未表示 DeepSeek V4 已在昇腾上完成全规模训练——本次开源证明的是「V4 训练算子在昇腾上有高性能实现」,而非「V4 已用昇腾训练完成」。
华为一方的配套动作包括:在 CANN 社区开源部署与训练实践,附 DeepSeek V4.1-Flash 每卡输出吞吐实测数据,为使用者提供从环境搭建到训练推理的参照路径。两侧开源内容合起来,构成「模型厂软件栈 + 芯片厂工具链」的完整组合。
三、数据与规格
| 项目 | 数据 |
|---|---|
| 公布时间 | 2026 年 9 月 30 日 |
| 开源组件 | TileLang 昇腾版、DeepGEMM-Ascend、TileKernels、FlashMLA、DeepSelect、DeepEP-Ascend、DeepJIT |
| 实测硬件 | 昇腾 950DT NPU |
| 实测软件栈 | CANN 9.2.0、Python 3.12、PyTorch 2.13.0、torch_npu 2.13.0rc1 |
| Dense GEMM 利用率 | 最高 99.8%(BF16)、99.5%(FP8)、98.3%(FP4) |
| DeepEP-Ascend 带宽 | EP8 下 Dispatch 373–375 GB/s、Combine 345–347 GB/s |
| Dispatch 带宽上限达成度 | EP32 及以下约 90%–95%(官方口径) |
| 商用固件 | 华为 Q3 商用 HDK(Atlas 850E),计划约 2026 年 10 月 15 日公开 |
| 开源协议 | DeepGEMM-Ascend、TileKernels、DeepJIT 等为 MIT 协议 |
| TileKernels 昇腾后端要求 | Ascend 950 NPU、CANN 9.2.0 及以上 |
四、行业观察
从「移植」到「同构」的开源策略。 过去国产芯片适配大模型训练框架的常见路径,是把英伟达生态中的某个库单独移植过来,接口与性能都难以对齐。本次发布的取向不同:以 DSL(TileLang)承载算子、以运行时自动选后端的方式,让同一套 Python API 覆盖两种硬件,昇腾与英伟达版本一一对应、接口对齐。对模型厂商而言,训练基础设施不再绑定单一硬件供应商;对芯片生态而言,借用的是一个已被 V4 训练验证过的软件栈,而非从零建设。
国产算力软件生态的关键补课。 大模型训练的实际性能瓶颈,除算力本身外,集中在计算切分、数据搬运、多卡协同等基础设施层面。英伟达 CUDA 生态的护城河很大程度由这类「脏活」组件构成;昇腾 CANN 工具链近年持续推进,而 DeepSeek 把自家生产级训练组件整体开源到昇腾平台,等于给这条软件栈补上了经过最大规模生产验证的一环。官方给出的密集可复现数据(bench_msprof 脚本、tests 目录、明确的测试配置)也让性能声明具备可验证性,区别于单纯的宣传口径。
仍需观察的三个问题。 其一,商用固件尚未发布:当前性能数据出自 PoC HDK 手工配置,公开用户预计 10 月 15 日前后拿到正式版本后才能独立复现;Combine 在大 EP 规模下与 URMA 的 HBM 争用也仍待华为固件改进。其二,全规模训练未证实:官方未表示 V4 已在昇腾上完成全规模训练,算子级适配与万卡级训练稳定性之间仍有距离。其三,版本覆盖:现有内核的验证范围集中在 Ascend 950 系列与 CANN 9.2.0,更早代际的昇腾设备与 CANN 版本尚未由这些实测数据覆盖。这三个问题决定了这套组件从「开源发布」走向「生产可用」的实际节奏。
举手提问