2026 年 9 月 30 日，DeepSeek 把一整套面向华为昇腾平台的底层算子栈开源了出来：注意力算子库 FlashMLA 发布昇腾算子，通用算子库 TileKernels 加入昇腾后端，计算库 DeepGEMM-Ascend 与通信库 DeepEP-Ascend 两个新仓直接开张，覆盖矩阵运算、分布式通信、通用算子与稀疏注意力四类核心能力。按新闻口径（AI工具集 09-30 期每日快讯，转述 DeepSeek 官方），这批组件与此前英伟达平台上的对应组件一一对应，其中 TileLang 昇腾版已支撑 DeepSeek V4 系列模型训练算子的高性能实现，多项测试接近硬件上限。本文逐仓拆开看，全文 star 数据均为 2026-10-01 GitHub API 快照。
这件事的背景值得先摆一下。在此之前，头部大模型公司里公开把训练线搬到国产芯片上的只有美团一条路（LongCat 2 的国产芯片训练线，见站内美团 LongCat 2 热点篇，该文事实口径独立成篇，本文以事实库为准）。美团讲的是怎么训，DeepSeek 这次讲的更底层：旗舰模型在昇腾上跑起来，算子层长什么样、性能贴到哪条线，全部摊开到源码级别。模型权重早就开源了（DeepSeek-V4.1 开源篇见站内DeepSeek V4.1 热点），这次连「权重是怎么在这块芯片上被计算和通信的」也一并交了出来。
四仓总览：一张表看定位
先把四个仓库摆进一张表。star 均为 2026-10-01 GitHub API 快照，许可证以仓库根目录 LICENSE 原文为准。
| 仓库 | star | 许可证 | 角色 |
|---|---|---|---|
| TileKernels | 1883 | MIT | TileLang DSL 通用算子库，NVIDIA 与昇腾双后端 |
| DeepGEMM-Ascend | 418 | MIT | 矩阵乘法计算库，DeepGEMM 的昇腾移植 |
| DeepEP-Ascend | 190 | 根目录暂无 LICENSE 文件 | MoE 专家并行通信库，DeepEP 的昇腾实现 |
| FlashMLA | 13015 | MIT | MLA 稀疏注意力算子库，同一仓覆盖双平台 |
四个仓库的分工可以这样理解：TileKernels 是最上层的通用算子库，用 TileLang 这门多后端领域专用语言写成；DeepGEMM-Ascend 扛矩阵乘法，这是大模型里最大头的算力消耗；DeepEP-Ascend 管 MoE 训练推理中最要命的专家并行 all-to-all 通信；FlashMLA 管注意力，推理时延的大头。四层拼起来，恰好是大模型训练推理中最吃硬件的几块。
TileKernels：1883 颗星，一套 API 跑两种芯片
TileKernels 是四个仓里最早创建的（2026-04-22），也是唯一走 pip 分发的（包名 tile-kernels），2026-09-30 推送了那次关键更新。它基于 tile-ai/tilelang 这门支持多硬件后端的领域专用语言，提供数十个深度优化算子：MoE Routing 的 top-k 专家选择与打分；量化，覆盖 per-token、per-block、per-channel 三种粒度的 FP8/FP4 cast 与反量化，还有融合 SwiGLU 加量化的算子；Engram 门控，融合 RMSNorm、前反向与权重梯度归约；流形超连接 mHC，含 Sinkhorn 归一化；外加 RoPE 和 Rand 算子。
两个信息最值得划线。其一，README 官方自报：大多数算子的性能接近硬件的计算吞吐或内存带宽上限，而且全部算子已用于内部训练与推理任务——不是演示仓，是生产仓。其二，2026-09-30 的 News：加入华为昇腾支持，沿 NVIDIA 路线新增了第二个后端，运行时自动选择，同一套 Python API 可以同时跑 NVIDIA GPU 与昇腾 NPU。不需要改代码切分支，后端在运行时自己挑。README 的致谢一节原文感谢了华为在 Tile Kernels 昇腾后端开发中的技术支持与工程投入。
新闻口径里还有一句值得注意：TileLang 昇腾版已支撑 DeepSeek V4 系列模型训练算子的高性能实现。也就是说这批算子不只在跑 benchmark，而是真的在给 V4 系列的训练供弹药。环境要求方面，Python 3.12 以上、PyTorch 2.13 以上、TileLang 0.1.15 以上；CUDA 后端要 SM90 或 SM100 架构加 CUDA 13.1 以上；昇腾后端要 Ascend 950 NPU 加 CANN 9.2.0 以上。
DeepGEMM-Ascend：418 颗星，API 完全兼容的矩阵乘法库
DeepGEMM-Ascend 是 DeepGEMM 在昇腾平台上的移植，2026-09-29 创建，09-30 首发（README News 原文：初始发布即支持 Ascend 950 设备）。它最实用的一个设计是完全兼容 DeepGEMM 的 API，连包名都沿用 deep_gemm：用户装上这个包，就能直接沿用其他平台上 DeepGEMM 的 API 和开发流程，迁移成本被压到接近零。功能上支持 BF16、FP8、FP4 三种 GEMM，外加 MQA logits 和 MegaMoE 算子。
工程实现上，它对昇腾的矩阵乘加原语（MAD）做了一层轻量抽象，把分形布局、对齐约束、地址计算这些繁琐细节藏起来，让 GEMM kernel 的实现保持简洁；同时大量采用昇腾特有的优化技术，比如稀疏数据加载和基于协程的流水线。README 官方自报：多种矩阵形状上接近硬件极限性能。注意这里没有给出具体 TFlops 数字，性能结论以仓库 README 为准。性能优化参考价值也被明确写在 README 里：这些实现可以作为昇腾平台极致性能优化的参考。
依赖项要先看清：CANN 9.20 工具链（提供 bin/bisheng 与 bin/ld.lld）、torch_npu、Python 3.10 以上、支持 C++20 format 特性的编译器，另外依赖 tilelang（HC prenorm kernel 用到）。还有一个容易踩的细节坑：README 原文明确，昇腾侧的 scaling factor 格式与英伟达不同——沿 K 维的一对 UE8M0 scaling factor 被打包进一个 int16，按 MN-major 顺序存储。API 兼容不等于数据格式无脑互通，从英伟达侧搬流水线时这一层要自己适配。
DeepEP-Ascend：190 颗星，最强的一块硬骨头，也是合规上最模糊的一块
DeepEP-Ascend 是 MoE 专家并行通信库，2026-09-30 创建并推送。训练、推理 prefill、decoding 三种场景共用同一套 EPBuffer 接口和 deep_ep 包名，实现遵循 DeepEP 的 EPBuffer V2.5 API，具体模式与流行为按昇腾特性实现。除 EP 之外，README 还提到流水线并行、面向 CP/DP 的 Bucket 集合通信以及 Engram 远端内存访问等通信原语，标注为开发中。依赖上，它靠 DeepJIT 子模块（deepseek-ai/DeepJIT，359 颗星，MIT）在运行时编译 Ascend C 内核。
验证栈写得非常具体：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；软件侧要有带 Ascend C、Bisheng 编译器、HCCL、HCOMM 的 CANN。性能测试条件（README）：16384 tokens/rank、hidden 7168、top-6 路由 256 专家、32 个 AI core 与 64 个 AIV、Dispatch 用 FP8 行主序 scale、combine 用 BF16、10 次预热加 50 次采样。README 以图表形式给出各 EP 规模下的带宽结果，具体吞吐数字以仓库原文为准，本文不转引。测试条件本身透露的信息量不小：hidden 7168 对应 V4 系列的模型宽度，top-6 加 256 专家对应 MoE 结构的经典配置，测试是按真实模型形态来设计的，不是随手挑的好看形状。
这个仓有一个必须单独拎出来说的问题：许可证缺位。截至 2026-10-01 快照，DeepEP-Ascend 仓库根目录没有 LICENSE 文件——GitHub API 的 license 字段为空，contents 接口列出的根目录文件里也没有 LICENSE。对照之下，同批的 DeepGEMM-Ascend 与 TileKernels 的 LICENSE 原文均为 MIT（已核对），FlashMLA 也是 MIT，唯独它缺位。没有许可证文件意味着默认版权保护：你可以看，但严格来说不能复制、修改、再分发。考虑到它又是四仓里工程上最硬、企业用户最关心的通信库，这个缺口对采用决策的影响不小。建议按合规待观察处理，等官方补上 LICENSE 再做生产引入判断。也正因为它是训练侧绕不开的一环——MoE 架构下专家并行的 all-to-all 通信效率直接决定扩展效率——这个空着的位置才格外扎眼。
FlashMLA：13015 颗星，410 TFlops 只算 95 分
FlashMLA 是四仓里星数最高、资历最老的（2025-02-21 创建），英伟达侧和昇腾侧共用这一个仓。2026-09-30 的 News：发布华为 Ascend 950 NPU 上的稀疏注意力 prefill 与 decoding kernel，按 README 官方口径，prefill 达到 410 TFlops，为硬件峰值的 95%；decoding 达到 360 TFlops，为硬件峰值的 83%。同时发布了一篇讲算法与优化技术的深度报告（仓库 docs 目录下 20260930-ascend-prefill-deep-dive.md，中英双版），把优化手法摊开讲，这在算子开源里算少见的诚意。
对照英伟达侧更有感觉：B200 加 CUDA 13.3 上，融合了 Q-norm、Q-RoPE、attention、O-RoPE 与 FP8 cast 的融合 kernel，prefill 至 1460 TFlops、decoding 950 TFlops。两边架构不同，数字不能直接横比，但昇腾侧 prefill 贴到 95% 峰值这一点说明算子质量是第一梯队。本次更新还将融合 kernel 的 decoding 性能优化了约 10% 到 15%（README 官方口径），不过要注意：这个融合 kernel 目前仅支持 CUDA，昇腾不支持，README 原文写得很明确。
FlashMLA 支撑 DeepSeek-V4.1 模型在 NVIDIA GPU 与华为昇腾 NPU 上的推理，功能覆盖 prefill、FP8/FP4 KV 缓存的 decoding、融合 attention 前后小操作的融合 kernel，以及稠密 attention 的 prefill 与反向算子。标题里那句「410 TFlops 只算 95 分」就是这个意思：prefill 已经贴近硬件上限，decoding 的 83% 留有余量，也是后续版本最值得盯的改进点。
与英伟达侧对应仓的关系：一一对应，牌摊在明处
新闻口径说这批组件与英伟达平台组件一一对应，落到仓库层面是这样：DeepGEMM-Ascend 对应英伟达侧的 DeepGEMM（7907 颗星，MIT，快照同日）；DeepEP-Ascend 对应 DeepEP（10231 颗星，MIT）；FlashMLA 则是同一个仓直接覆盖双平台（13015 颗星，MIT）。昇腾侧两个独立新仓 star 体量小很正常——创建日期是 09-29 和 09-30，它们刚出生几天；TileKernels 本身就是双平台仓，1883 颗星已经不算低。此外还有两个背景仓顺带一提：DeepJIT（359 颗星，MIT）是 DeepEP-Ascend 的子模块，clangd-ascend（14 颗星）是 IDE 支持，许可证字段标注为 NOASSERTION。
把这件事放进更大的图景里看：模型权重开源（V4.1）、训练推理算子开源（这一批）、双平台后端开源（TileKernels 与 FlashMLA），DeepSeek 是目前把整条技术栈摊得最开的头部团队。对昇腾生态，这批仓库提供了旗舰模型级的参考实现——怎么在 MAD 原语上抽象 GEMM、怎么用 HCCL/HCOMM 与 URMA 做 EP 通信、怎么把稀疏注意力贴到 95% 峰值，全是可读可改的代码。对行业，它把国产芯片软件生态的成熟度从「能跑」推进到了「有公开算子栈可对账」。训练基建这条线，站内还有Agentic RL 框架横评与MiMo 加 verl 实战篇：那两篇讲的是框架层怎么组织训练，本文讲的是框架脚下的算子层，正好互补。
短板也要说透。第一，昇腾侧算子栈验证环境门槛不低：DeepEP-Ascend 要 950 系列加 UBMEM 连通与特定 CANN 版本，DeepGEMM-Ascend 绑定 CANN 9.20，普通开发者手上没有卡就只能围观。第二，FlashMLA 的融合 kernel 等部分能力仍仅支持 CUDA，昇腾侧覆盖还在逐块补齐。第三，DeepEP-Ascend 的 LICENSE 缺口如前所述，商用前必须等它补齐。
写在最后
四仓摆在一起，DeepSeek 这次开源的信号大于数字本身：算子层不再是黑箱，昇腾平台第一次拥有了旗舰模型同款的公开实现，性能口径也全部给到了硬件峰值百分比这种可核对的写法。410 TFlops 只算 95 分，剩下那 5 分和 decoding 那 17 个百分点，就是这套栈明牌写着的下一步。想跟进昇腾训练线的团队，建议从 TileKernels 的 pip 包和 FlashMLA 的深度报告入手，DeepGEMM-Ascend 适合有算子开发经验的团队精读，DeepEP-Ascend 先观望许可证。
常见问题
Q1： 四个仓库分别管什么，都要装吗？ A1： TileKernels 管通用算子（MoE 路由、量化、Engram 门控、mHC、RoPE 等），DeepGEMM-Ascend 管矩阵乘法，DeepEP-Ascend 管 MoE 专家并行通信，FlashMLA 管稀疏注意力。按需安装：做推理优先 FlashMLA 加 DeepGEMM-Ascend，做训练再加 DeepEP-Ascend，通用算子需求装 tile-kernels 即可。
Q2： DeepGEMM-Ascend 与 DeepGEMM 完全 API 兼容，代码能直接搬过去吗？ A2： 接口层面可以，包名同为 deep_gemm，沿用原开发流程。但 README 原文明确，昇腾侧 scaling factor 格式与英伟达不同：沿 K 维的一对 UE8M0 值打包进 int16 并按 MN-major 存储。涉及量化数据的流水线要自己适配这一层，环境上还需 CANN 9.20、torch_npu 与 Python 3.10 以上。
Q3： DeepEP-Ascend 没有 LICENSE 文件，能直接商用吗？ A3： 不能想当然。截至 2026-10-01 快照，该仓库根目录没有 LICENSE 文件，GitHub API 的 license 字段为空，默认版权保护下严格来说只能查看不能复制、修改、再分发。同批其他三仓均为 MIT。建议按合规待观察处理，等官方补上许可证文件再做生产引入。
Q4： FlashMLA 昇腾版 410 TFlops 是什么水平？ A4： 这是 README 官方口径：Ascend 950 NPU 上稀疏注意力 prefill 410 TFlops，为硬件峰值 95%；decoding 360 TFlops，为峰值 83%。prefill 已贴近硬件上限，decoding 还有余量。对照参考：B200 上融合 kernel prefill 至 1460 TFlops、decoding 950 TFlops，架构不同不可横比。
Q5： 想上手这批仓库需要什么环境？ A5： 硬件层面需要 Ascend 950 系列 NPU；软件层面普遍要求 CANN 9.2.0 以上（DeepGEMM-Ascend 为 CANN 9.20）与 torch_npu，TileKernels 另需 Python 3.12 加 PyTorch 2.13 加 TileLang 0.1.15，DeepEP-Ascend 的验证栈是 Python 3.12 加 PyTorch 2.13.0+cpu 加 torch_npu 2.13.0rc1 并要求偶数 rank 与 UBMEM 连通。没有昇腾设备的话，先读 FlashMLA 深度报告和各仓 README 最省时间。