2026 年 9 月 30 日，DeepSeek 和华为开源了一套面向华为昇腾芯片的软件组件，包含 6 个仓库，覆盖了矩阵乘法，注意力，通用算子库，Top-k 选取，分布式通信和编译工具。华为还发布了 128 卡昇腾 950 的超节点方案 SuperPoD Flex，这些组件就是为这个硬件做的。这些都不是 demo 性质的玩具，是大模型训练和推理里最耗费算力、最难优化的算子。官方称它们属于编译工具，高性能计算库和分布式通信库。没有调度层模块。
团队评估第二平台，手里就两把尺：一把量省力，一把量性能，跑到最终训练的团队对第二把最敏感。过去所有做跨平台的人问的都是同一个问题：怎么修改一套代码，让它同时在 NVIDIA 和别的芯片上跑起来。这个问法用了很多年，始终没有好答案。DeepSeek 换问的是：两颗芯片之间，哪些部分必须一样，哪些部分允许不一样。拆下来是四个小问题：模型代码要不要动，显存里的数据怎么摆，每个算子谁来写，开发工具共不共用。这批组件对每个小问题都交了答案，答案全部公开可查。
旧问题为什么始终没有好答案，得看一个物理事实。同样的矩阵乘法，同样的数字，存进显存的字节却完全不同：你调用的函数和参数都一样，两颗芯片摆这些数据的方式就是不一样。硬要在字节这一层统一，总有一颗芯片要多付数据转换的开销，甚至用不上自己最核心的计算部件。峰值要贴着硬件跑，贴着硬件就没有统一，这个矛盾在物理层就成立，跟软件写得好不好没有关系。
全行业跟这个矛盾较劲了多年。Google、OpenAI 和摩尔线程各选过一层去统一，在自家地盘都成功，换到第二平台都撞了墙。用 AI 辅助写代码变成现实之后，账面起了变化：第二套实现的人工成本被打下来，芯片厂商、框架组、模型组重新都动了起来。不过成本降下来只解决了愿不愿意写；四个小问题里每一项锁什么、放什么，公开场合从来没有人给出过一份完整答案。
这批组件把答案摆上了公开桌面，同等范围的开源先例，本次调研尚未找到。真实生产算子换到新硬件，哪些可以原样复用，哪些必须重写，重写要花多少工程代价，都有了可以逐行查对的代码依据。下文先看三条旧路线各自卡在哪，再看物理约束怎么压天花板，接着看 AI 写代码改变了什么、没改变什么，然后拆这批组件怎么把兼容按层标价，最后算这套做法的代价。
先看三条旧路线，看它们各自在两难里选了什么位置。过去十年，为了摆脱对单一芯片厂商的依赖，社区探索了三种典型路线：Google 把希望放在编译器上，OpenAI 放在写算子的语言上，芯片厂商放在API兼容上。三条路都在硬件性能和易用性之间选了不同的平衡点。
XLA/JAX：用计算图编译器签兼容合同。画图就行，细节交给编译器。在自家定制芯片上软硬结合跑得快，但在通用 GPU 上，尤其是对大模型关键的高吞吐矩阵乘法和稀疏注意力，编译器生成的代码性能始终不如工程师手工优化，生产线上最热的算子还是靠专家手写。
第二条路是 OpenAI 主导的编程模型路线，典型工具是 Triton。它把契约签在算子编程模型这一层，让开发者用类 Python 的语法来写高性能 kernel。Triton 提供的抽象基本上和 NVIDIA 的硬件是一一对应的。这套抽象在它的原生硬件上几乎是免费的，但一旦挪到其他芯片架构上，那些看似理所当然的硬件假设就成了沉重的适配负担，第二平台的算子生态无论在成熟度上还是在调优深度上，都迟迟没能跟上主平台。
第三条路是芯片厂商常用的 API 克隆路线，比如摩尔线程的 MUSA。它把合同直接开在成熟生态的 API 表面，目标是让现成代码不用改就能在新芯片上编译运行。对着急跑起存量业务的客户来说，这种迁移门槛确实最低；但新芯片的流水线和微架构为了迎合旧接口的既有行为，往往处处受制，很难放开手脚发挥内部硬件的原生潜能。
这三条路线锁住的是不同层面的东西：
| 路线 | 谁在做 | 锁的兼容单位 | 卡在哪 |
|---|---|---|---|
| XLA / JAX | Google，框架方 | 锁计算图，kernel 交给编译器生成 | 最难的 LLM 算子（FP8 GEMM、稀疏 attention）上编译器生成的 kernel 到不了手写水平，NVIDIA 上的生产热点仍是手写路线 |
| Triton | OpenAI，框架团队 | 锁与 NVIDIA 硬件一一对应的编程模型 | 这把锁在 NVIDIA 上免费，换到别的芯片变成成本，第二平台的 kernel 生态没有接近 NVIDIA 侧 |
| MUSA | 摩尔线程，芯片厂商 | 锁 CUDA API 表面，让既有 CUDA 代码原样运行 | 新芯片只能模仿 CUDA 的行为，没有使用自家微架构的自由 |
| 边界推演 | 模型实验室 | 锁运算语义，最小兼容单位 | 各条路线锁的单位不同，单位越大，新硬件上让出的自由度越多。现在把运算语义锁起来的是模型实验室，因为只有它同时拥有算法和调用点；同等范围的开源先例，本次调研尚未找到。 |
三条旧路线都撞了墙，撞墙的方式说明同一件事：兼容签在哪一层，直接决定第二块芯片还剩多少自由。这件事为什么是真的，得看大模型本身的负载。训练和推理的吞吐瓶颈几乎全部集中在矩阵乘法、稀疏注意力和跨节点专家通信这三类算子上，它们必须把芯片的计算单元与带宽压榨到物理极限，容不得半点浪费。
具体地说，GPU 和 Ascend 的硬件设计哲学不同：GPU 有独立的高速互联结构，让外部内存和内部缓存之间的数据传输可以和计算流水线无缝重叠。Ascend 把矩阵单元和向量单元做成两个独立核，片内搬运靠编程显式调度和同步。因此，指望从微架构指令里直接抹平二者差异，最终都会在流水线调度上翻车。
一个具体的例子是，在流行的 FP8 低精度矩阵乘法里面，为了保证精度，需要引入缩放因子。NVIDIA 的硬件要求把 4 个缩放因子塞进一个 32 位的整数中，然后按照一定的转置格式存在显存中，方便数据搬运部件读取；但是昇腾的硬件规则要求我们把 2 个缩放因子塞进 16 位的整数，对行和列的要求也完全不同。
如果非要把两边数据在字节层面的排布完全绑定，新芯片就会被迫花大量时钟周期去解包、转置、搬移数据，甚至动不了核心高速计算模块。合同定的越细，新芯片模仿旧芯片行为的代价越大，能达到的吞吐上限就越低。这种损失是局部代码调整根本弥补不了的，因为芯片内部的数据通路宽度和排布规则早就刻在硅片上了。
客观地说，XLA 在 TPU 上的成功和 Triton 在它自己核心平台上的高效率都证明了这些方法在它们的原生化领域是非常有效的。问题出在想实现同一个运算在两个不同的芯片上都能达到硬件峰值。这是一个物理上很难的问题。无论软件多么廉价，都无法改变硬件如何组织数据。
物理层面打不开的僵局，AI 编程从成本这一侧把它撬松了：写第二套实现的人工变便宜，这件事在这批仓库里能直接看到。编译工具 tilelang-ascend 的仓库里留着大量 AI 参与开发的工程痕迹：21 个供智能体调用的技能文件，4 个 agent 串成一条三阶段状态机，覆盖算子设计、实现和性能调优；CI 会自动拉取示例跑基准测试，提交记录里频繁出现模型协同签名的记录。华为的开发者社区里也公开给出过面向智能体写算子的规范，官方提过智能体半小时写完一个算子。这个数字只算写代码这一段，数值验收、跨版本回归、通信正确性这些兜底仍然是资深工程师的活。
拿两条老路推演一下，就能看出 AI 的边界在哪。继续走 MUSA 那条 API 兼容的路：AI 把兼容层写得再快，写出来的仍然是模仿，好用的问题解决了，性能天花板一步没动。想直接对齐别家芯片的底层指令：那条跑道每代都在变，对手换一个计算部件，写好的适配就全部作废。AI 能把代码写快，但前提是那里已经有约定；约定开在哪一层，还是得人来定。
所以问题就从有没有人力做两套实现，变成了合同签在哪一层。写代码变便宜之后，团队第一次有余裕认真算另一笔账：上层保持统一，底层按芯片各写一套，维护账和商业账能不能算得过来。
DeepSeek 对这笔账的回答，是接下来这批代码。先看 6 个仓库是什么：TileLang-Ascend（一套用 Python 写 kernel 的 DSL）、DeepGEMM-Ascend、TileKernels、FlashMLA、DeepSelect、DeepEP-Ascend。官方通报把它们定位为编译工具、计算库和通信库，没有调度层模块；在线 serving 调度与集群负载均衡是显式排除的测试条件，不在开源范围里。
这套软件处理两颗芯片差别的方式，直觉上就一句话：对上的承诺给足，对下的差别放开。模型怎么调用这些算子，两边接口完全一致，模型代码一行不用改；接口往下，两颗芯片各按各的方式干活。拆成四层：
| 层 | 策略 | 证据 |
|---|---|---|
| 调用，Python API | 锁死：同名同参 | 模型侧调用代码零改动 |
| 数据字节，缩放因子布局等 | 故意不同：保含义不保字节 | FP8 缩放因子：NVIDIA 4 个进 32 位整数，昇腾 2 个进 16 位整数，行列方向要求不同，用户侧转换函数处理 |
| kernel 实现 | 两套分写：同一门高级领域专用语言，_cuda.py/_asc.py
两份手写 |
engram_fused_weight：CUDA 版线程绑定，昇腾版核分工与缓冲区多版本 |
| 开发工具，DeepJIT | 真共享：一个库两边用 | NVIDIA 版 DeepGEMM/DeepEP 从 V2.5 起也换用它 |
四层怎么落地，看仓库代码里的三个地方。
第一处是数据在显存里怎么摆。同一个矩阵乘法，同样的数字，两边的字节排布不一样，前面说过的缩放因子打包就是例子。DeepSeek 没有把这个差异抹平，反而把它当成设计起点：格式转换集中到一个用户侧函数里，数据进门时先摆正格式，再交给算子。数学含义不变，字节怎么排由每种硬件自己做主。
第二处是每个算子怎么写。同一个算子备两份实现，一份给 NVIDIA，一份给昇腾，用同一门语言分别写成。拿仓库里最简单的算子 engram_fused_weight 对照，连线程怎么分工、数据放哪、怎么搬，两边没有一项相同。写成两份，是因为两颗芯片的执行方式差异大到没办法用一套代码兼顾。
第三处最说明问题：整套代码里唯一跨平台共用的，是开发设施。运行时编译、哈希缓存、加载这套工程环节抽成了一个单独的小库 DeepJIT，两边共用；连 NVIDIA 自己的矩阵乘法库和通信库，从 V2.5 起也换用了它。算子各写各的，造算子的工具只用一份。通信库是同一个思路：上层接口对齐，底层换成华为自己的通信栈 HCCL；没做完的功能在文档里公开列着，代码里是占位符。四层分工连起来看：算子语义免费锁死，字节差异付一个转换函数，峰值各付一套定制实现，工具层共享反而省一份。
这个做法的代价很直接。这批组件里最有分量的三个，矩阵乘法库的核心主体、注意力库 FlashMLA、Top-k 选取 DeepSelect，计算主体没有用任何高级语言，全部直接用 Ascend C 写底层实现；这里的直接，指的是不经过 DSL 生成，跟用不用 AI 辅助是两回事。Ascend C 是华为给昇腾芯片的底层编程语言，地位类似 CUDA C。其中 Top-k 的实现压在一个文件里，54KB、约 1800 行。换一块芯片，就意味着再写一遍这些代码。
成本付了这么多，收益到底有多大？这批组件公开了一批自测数字。先说结论：数字看着漂亮，但全部来自厂商自测，每个都带着一串先决条件，而且至今没有第三方验证。下面一个个看。利用率最高的数字来自矩阵乘法库。DeepGEMM-Ascend 发布的自测表格中，950DT 上一个 BF16 dense GEMM case 报告 431 TFLOPS（M=4096, N=7168, K=16384，CANN 9.20，cold L2）；以维护者引用的 432 TFLOPS 规格为分母，约为 99.8%。本次核查未获得独立复现，GitHub issue #1 对规格口径提出了质疑。该测试对象是直接用底层语言写成的原生 Ascend C 内核，不经过 DSL，所以不能拿它证明 DSL 也能跑到芯片峰值。同时在 GitHub issue #1 讨论区中，有开发者指出若按照特定高规格档位的浮点算力为分母，该算子的实际利用率约为 85% 左右，issue 至今 open。最昂贵的部分没有被 AI 变便宜。注意力算子上，FlashMLA 9/30 版的 prefill 报告 410 TFlops（硬件峰值 95%）、decode 报告 360（83%），V4.1 only，dav-3510，厂商自测。
通信数字的先决条件最重。DeepEP-Ascend 的带宽数据来自提供给 DeepSeek 的 PoC HDK，一套内部试验专用的硬件开发套件，与额外人工配置，该配置并未公开分发；README 将计划在 2026 年 10 月中旬公开的商用 HDK 列为后续部署基线，该日期是厂商发布计划。带宽的完整口径是：DeepEP-Ascend 在 950DT、指定 PoC 配置的自测中，EP8 的 FP8 dispatch 报告 373–375 GB/s，BF16 combine 报告 345–347 GB/s；测试使用 netlayer 1，即超节点外部 Clos 网络；计时包含 issue 和 drain，不包含 final epilogue。这是超节点外部网络层的数字，节点内部的高速互联是另一回事。
端到端推理吞吐是另一个口径：华为公布，在其 EP32、128K context、Dspark 投机接受率 0.85 的离线测试设置下，V4.1-Flash 报告 2469 tok/s/card（TPOT 5 ms）；测试明确排除了 serving 调度与框架负载均衡影响，本次核查未获得独立复现。在相同的离线测试环境下，对应于 TPOT 10 ms 的吞吐数据报告为 5102 tok/s/card。
这些数字之外还有一个朴素的事实：至今没有第三方验证。代码公开后几天里，外部开发者提交的所有 PR 都由作者自己标注了“未在 NPU 上实际运行”；无公开证据表明，有第三方团队在独立的 950 集群上完整跑通了这套软件栈。仓库的 issue 区也记着早期的粗糙：通信库从开源开始就没有 LICENSE 文件，40 分钟内就有人提出，该问题未解决(Issue #1)；矩阵乘法库深度绑定 950 独有的机制，vLLM-Ascend 团队实测在上一代 Atlas A3 上跑不通(Issue #3)；Top-k 组件还有第三方对算法原创性和引用提出优先权异议(Issue #13)。训练方面，官方通报提到训练用到的算子在昇腾上都有对应的高性能实现，同时写明“TileLang 路线首先在英伟达成熟平台上得到了验证”；到目前为止，无公开证据表明相关模型已经在昇腾集群上完成端到端的全量预训练。
代价、数字、空白都看完了，回到开头的问题。这批组件真正交出去的，是一份可以逐行核查的边界：哪些必须兼容，哪些放开，各自值多少。DeepSeek 的昇腾组件展示了一条务实路线：保留模型运算的接口与语义，共享开发设施，允许关键 kernel 按硬件重写；它公开的是分层复用的工程边界，不是已经完成的通用性能可移植性突破。摆脱 CUDA 依赖的关键，是决定哪些东西不必兼容。