2026 年 9 月 30 日，DeepSeek 做了一件此前没有第二家大模型公司做过的事：把自己赖以训练和推理的整套基础设施算子栈，从英伟达平台平移到华为昇腾，并在一天之内全部开源。这一天里陆续亮相的，包括基于 TileLang 的深度优化算子库 TileKernels、矩阵计算库 DeepGEMM 的昇腾移植版 DeepGEMM-Ascend、专家并行通信库 DeepEP 的昇腾实现 DeepEP-Ascend，以及 13015 星老仓库 FlashMLA 新增的昇腾注意力算子。
官方口径（经 AI工具集每日快讯转述）是这样描述的：这批组件与此前英伟达平台上的组件一一对应，覆盖矩阵运算、分布式通信、稀疏注意力等核心能力，多项测试接近硬件上限。其中 TileLang 的昇腾版已经支撑 DeepSeek V4 系列模型训练算子的高性能实现。最抓眼球的数字来自 FlashMLA：在华为 Ascend 950 NPU 上，稀疏注意力 prefill 达到 410 TFlops，相当于硬件峰值的 95%；decoding 达到 360 TFlops，相当于硬件峰值的 83%（README 官方口径）。
需要强调的是，「多项测试接近硬件上限」这一说法是官方口径，来自官方转述渠道，而非第三方评测结论。这两个数字意味着什么？对国产算力生态又意味着什么？这篇文章把四件套逐一拆开看，也把光环之下的几个问号摆上桌面。
一天四件套：一张英伟达栈对照表
先说清楚这四件套在整套训练推理栈里各管一段。大模型的训练与推理可以粗分为三层：最底层是算子，也就是矩阵乘、注意力、量化这些最基本的计算单元；中间是通信，专家并行架构下不同设备之间要高频交换数据；最上层才是模型与训练框架本身。DeepSeek 此前在英伟达侧已经陆续开源了 DeepGEMM、DeepEP、FlashMLA 等基础设施组件（均采用 MIT 许可证），这次的昇腾侧组件几乎可以逐一对上号：
| 仓库 | 在整套栈中的角色 | 星标数 | 许可证 |
|---|---|---|---|
| TileKernels | 基于 TileLang 的算子库，9 月 30 日起新增昇腾后端 | 1883 | MIT |
| DeepGEMM-Ascend | DeepGEMM（英伟达侧 7907 星）的昇腾移植 | 418 | MIT |
| DeepEP-Ascend | DeepEP（英伟达侧 10231 星）的昇腾实现 | 190 | 根目录暂无 LICENSE 文件 |
| FlashMLA | 13015 星的注意力算子库，9 月 30 日加入昇腾算子 | 13015 | MIT |
（星标数为 2026 年 10 月 1 日快照数据。此外还有一个背景仓 deepseek-ai/clangd-ascend，14 星，许可证状态为 NOASSERTION，一句话带过即可。）
四件套里最先值得展开的是 TileKernels。这个仓库 2026 年 4 月 22 日创建，原本是一套基于 TileLang 的深度优化算子库——TileLang 本身是 tile-ai/tilelang 这个多后端领域专用语言，开发者用类似分块（tile）的写法描述计算，再由编译器生成不同硬件上的高效实现。DeepSeek 在其上构建了几十个算子：MoE Routing 的 top-k 专家选择、量化方向的 per-token、per-block、per-channel 三种粒度的 FP8 与 FP4 cast 及反量化（含融合 SwiGLU 与量化）、Engram 门控（融合 RMSNorm、前反向与权重梯度归约）、流形超连接 mHC（Sinkhorn 归一化），以及 RoPE、Rand 等基础算子。官方说明里有两个关键表述：大多数算子接近硬件计算吞吐或带宽上限；全部算子已用于内部训练与推理。
9 月 30 日的新闻是：TileKernels 沿 NVIDIA 路线新增了第二后端，也就是华为昇腾，运行时自动选择后端，同一套 Python API 可以同时跑 NVIDIA GPU 与昇腾 NPU。环境要求上，CUDA 后端需要 SM90 或 SM100 架构与 CUDA 13.1 及以上版本，昇腾后端需要 Ascend 950 NPU 与 CANN 9.2.0 及以上版本，整体要求 Python 3.12+、PyTorch 2.13+、TileLang 0.1.15+，pip 包名为 tile-kernels。值得一提的是仓库致谢部分的原文：感谢华为在 Tile Kernels 昇腾后端开发中的技术支持与工程投入。这句话侧面说明这不是 DeepSeek 单方面的移植，而是有华为工程团队深度参与的联合工作。
第二件是 DeepGEMM-Ascend。矩阵乘（GEMM）是大模型里占比最高的计算，Transformer 的每一层都压在它上面，GEMM 跑不满硬件，整卡算力就都在漏。它是 DeepGEMM 的昇腾移植版，与 DeepGEMM 完全 API 兼容，连包名都保持 deep_gemm 不变，支持 BF16、FP8、FP4 三种精度的 GEMM，以及 MQA logits、MegaMoE 算子。技术实现上有两处昇腾特色：一是对昇腾 MAD（矩阵乘加）原语做了轻量抽象，把分形布局、对齐约束、地址计算这些底层细节藏起来；二是采用稀疏数据加载、协程流水线等昇腾特有的优化手段。2026 年 9 月 30 日的首发版本支持 Ascend 950 设备，运行环境要求 CANN 9.20（依赖 bin/bisheng 与 bin/ld.lld）、torch_npu、Python 3.10 及以上、支持 C++20 format 特性的编译器，并依赖 tilelang 提供 HC prenorm kernel。README 官方自报在多种矩阵形状上接近硬件极限性能——注意，这里没有给出具体 TFlops 数字，我们也就不替它编。还有一个容易被忽略的工程细节：昇腾侧的 scaling factor 格式与英伟达不同，UE8M0 沿 K 维打包进 int16、以 MN-major 方式存储（README 原文）。这意味着「平移」不是复制粘贴，量化数据排布都要按昇腾硬件重新设计。
第三件是 DeepEP-Ascend。DeepSeek 的 MoE 架构把专家分布在不同设备上，token 要跨设备路由，通信效率直接决定集群利用率，这就是 DeepEP 这类专家并行通信库存在的意义。它实现并遵循 DeepEP 的 EPBuffer V2.5 API，训练、推理 prefill、decoding 三种场景共用同一 EPBuffer 接口与 deep_ep 包名，模式与流行为则按昇腾特性实现。官方给出的验证栈相当具体：Ascend 950DT、CANN 9.2.0、Python 3.12、PyTorch 2.13.0+cpu、torch_npu 2.13.0rc1，外加 PoC HDK 配置；硬件层面要求 Ascend 950 (A5) 的 UBMEM 连通性与 UBC_CTP、URMA 通道，且要求偶数 rank；软件层面依赖 CANN 家族的 Ascend C、Bisheng 编译器、HCCL 与 HCOMM。README 中给出的性能测试条件是 16384 tokens 每 rank、hidden 7168、256 专家中 top-6 路由、32 个 AI core 与 64 个 AIV、Dispatch 走 FP8 行主序 scale、combine 走 BF16，10 次 warmup 加 50 次采样——但具体吞吐数字只出现在图表里，没有文本数字，所以本文不引用任何具体吞吐值。这个仓库还依赖 DeepJIT 子模块（deepseek-ai/DeepJIT，359 星，MIT）。
410 TFlops 与 95% 峰值：先看官方口径
四件套中数据最完整的是 FlashMLA。DeepSeek 系列模型用的是 MLA（多头潜在注意力）加稀疏注意力组合，推理时这部分计算最吃算力，也是各硬件厂商最头疼的优化对象。这个 2025 年 2 月 21 日创建的仓库本就是 DeepSeek 开源矩阵里的明星（13015 星，MIT），9 月 30 日发布昇腾注意力算子：面向华为 Ascend 950 NPU 的稀疏注意力 prefill 与 decoding kernel，prefill 达 410 TFlops，为 95% 硬件峰值；decoding 达 360 TFlops，为 83% 硬件峰值（README 官方口径）。仓库同时发布了算子算法与优化技术的深度报告，位于 docs/20260930-ascend-prefill-deep-dive.md，中英双版本，想看技术细节的读者可以直接去读原文。
对照英伟达侧：在 B200 加 CUDA 13.3 环境下，融合 norm-RoPE-attn-RoPE-cast kernel 的 prefill 至 1460 TFlops、decoding 950 TFlops。这里必须提醒两点。第一，410 与 1460 不可直接横比：硬件不同，kernel 形态也不同——1460 那个数字恰恰来自融合 kernel，而昇腾侧目前并没有同款融合 kernel。第二，95% 与 83% 都是 README 官方自报口径，第三方复现结果还需要时间。另外 README 提到，围绕昇腾的优化工作让 decoding 融合算子提速约 10% 到 15%（官方口径），同时明确写出：该融合 kernel 目前仅支持 CUDA，昇腾不支持。
从功能面看，FlashMLA 现在支撑 DeepSeek-V4.1 模型在 NVIDIA GPU 与华为昇腾 NPU 上的推理（关于 V4.1 模型本身的开源细节，见本站此前文章：DeepSeek V4.1 Flash 开源），支持 prefill 与 FP8、FP4 KV 缓存的 decoding，支持融合 attention 前后小操作的融合 kernel，以及稠密 attention 的 prefill 与反向算子。换言之，稀疏注意力这条 DeepSeek 模型架构里最吃算力的路径，如今在两家硬件上都有了对应的开源实现。
对国产算力生态意味着什么
把视野拉远一点，这件事的分量才显出来。在公开层面，此前只有美团把 LongCat 2 的训练线建在国产芯片上，并写了完整的技术复盘（见本站文章：美团 LongCat 2 如何把训练线搬上国产芯片）。美团证明的是「一家公司可以在国产芯片上训练大模型」，而 DeepSeek 这次的动作不同：它没有只交出一个训练好的模型或一份复现报告，而是把整套算子栈连同工程细节全部开源，让整个行业都能在这个基础上继续盖楼。
「一一对应」的价值也在这里。对算子库来说，跨硬件移植最难的不是写出能跑的版本，而是把性能磨到接近硬件上限——这需要同时吃透算法和硬件微架构，往往要数月的高强度调优。TileKernels 的多后端设计意味着同一套 Python API 在两家硬件上行为一致；DeepGEMM-Ascend 连包名都不改，存量代码理论上迁移成本极低；DeepEP-Ascend 遵循同一套 EPBuffer V2.5 API。对其他想做国产芯片适配的团队来说，这是一份可以直接对照阅读的「移植教科书」——哪些算子要重写、哪些数据格式要换（比如前述 UE8M0 的 scaling factor 排布）、哪些通信通道要依赖昇腾的 UBMEM 与 UBC_CTP、URMA，答案都摊在仓库里。
再往深一层看，TileLang 作为多后端 DSL 正在成为这类移植工作的底座，TileLang 昇腾版支撑 DeepSeek V4 系列训练算子的官方表述说明它已经过了实战检验。本站在开源旗舰模型对比一文中讨论过开源生态的竞争格局，这次的四件套补上了开源生态里最缺的一块：不是模型权重，而是让权重跑起来的基础设施。对国产算力来说，缺算子、缺工具链、缺工程范式的老问题，被头部模型公司用开源的方式补上了一角。
冷思考：光环之下的三个问号
热度之外，有三个问号值得冷静摆出来。
第一，DeepEP-Ascend 截至 2026 年 10 月 1 日快照，仓库根目录没有 LICENSE 文件，GitHub API 的 license 字段为空。对比之下，DeepGEMM-Ascend 与 TileKernels 的 LICENSE 原文均为 MIT。这意味着想引入 DeepEP-Ascend 的团队暂时处在合规待观察状态——没有许可证的代码在默认版权规则下不可自由使用，只能观望等官方补齐。对一个高调开源日而言，这是最显眼的一处缺口。
第二，融合 kernel 仅支持 CUDA。README 原文明确：融合 attention 前后小操作的融合 kernel 目前不支持昇腾。而英伟达侧 1460 TFlops 的 prefill 数字恰恰来自这个融合 kernel。换句话说，昇腾侧目前公开的是非融合路径的 410 TFlops，融合优化带来的额外收益（decoding 融合算子约 10% 到 15% 的提速）还停留在 CUDA 侧，昇腾侧何时补齐没有时间表。
第三，验证栈集中在 Ascend 950 系列。DeepGEMM-Ascend 首发支持 Ascend 950 设备，DeepEP-Ascend 的验证环境是 Ascend 950DT 加 CANN 9.2.0，TileKernels 的昇腾后端要求 Ascend 950 NPU 加 CANN 9.2.0 及以上。三份仓库的公开验证没有覆盖 950 之前的昇腾型号，存量昇腾设备能否跑起来、跑得如何，目前没有公开信息。再加上 DeepEP-Ascend 的性能数据只有图表没有文本数字，社区复现也需要等有 950 系列硬件的团队动手。
写在最后
一天开源四件套、与英伟达栈一一对应、昇腾侧注意力算子冲到 95% 硬件峰值（官方口径）——2026 年 9 月 30 日这一天，DeepSeek 把「国产算力能不能跑大模型基础设施」这个问题从争论变成了工程事实。但许可证缺口、融合 kernel 缺位、验证栈偏窄这三个问号也同样真实。接下来值得盯三件事：DeepEP-Ascend 何时补上 LICENSE，昇腾侧融合 kernel 何时跟进，以及拿到 Ascend 950 系列硬件的团队能否复现那些接近硬件上限的数字。开源的真正意义，是把一个公司在国产芯片上做训练推理的工程经验变成公共资产——这一次，DeepSeek 开的不只是代码，还有一个此前无人公开过的完整范式。