DeepSeek 开源算子工具大礼包，联手华为昇腾，手撕 CUDA 绑定！|cuda|deepseek|昇腾|知名企业|算子|算法|编译器

作者丨高允毅

编辑丨岑峰

刚刚，DeepSeek做了一个开源壮举：把给英伟达 GPU 写代码的全套工具链，原样铺到了华为昇腾 950 NPU 上。最核心的算子开发产品 TileLang 官宣原生支持昇腾。不仅如此，连同 DeepGEMM、DeepEP、FlashMLA、TileKernels、DeepSelect 五大核心计算与通信库也同步推出昇腾版本，直接对标英伟达平台的全套工具链。这意味着国产算力第一次在核心算子工具层面，做到了与英伟达生态的能力对标

在此之前，全球 AI 行业苦 CUDA 垄断久矣；而现在，DeepSeek 用一行行开源代码，为国产算力自主生态蹚出了一条康庄大道。

写算子的 “手工作坊” 时代，该结束了
做过 AI 底层开发的都知道，“写算子”一直是门槛最高、耗时最久的工作之一。一张芯片的理论算力上限很高，但如果底层算子写得不达标，芯片内部的计算单元就会有大量时间等数据流转，导致几万块钱的芯片可能只能发挥出 20% 的性能。写算子的目的，就是为了把这 20% 的硬件利用率提高到 95% 以上。拿最常见的矩阵乘法（GEMM）来说，要想榨干一张 AI 芯片的算力，开发者要手动管理三级缓存、调度线程块、搞定数据预取流水线，甚至要深入到寄存器分块和指令重排层面。为了优化上千行底层代码，资深工程师打磨好几周都是常态。这就导致了一个行业的核心冲突：模型算法迭代太快，芯片厂商官方算子库永远跟不上。过去行业里有三条路可选，但各有各的局限：第一条路是自己手写 CUDA：性能天花板最高，但开发效率极低，深度绑定英伟达生态，无法跨平台移植；第二条路是使用厂商的预制算子库：比如英伟达 cuBLAS，标准算子开箱即用、性能拉满，但灵活性极差，开发者无法修改内部逻辑、也无法做算子融合，新算法很容易卡在算力层；第三条路是使用高层 DSL 编译器：比如 OpenAI 的 Triton，大幅降低了开发门槛，但性能上限受限于编译器本身的优化策略，对昇腾、AMD 等非英伟达硬件原生支持度低，异构适配成本很高。而TileLang开辟了第四条路，终结算子开发的 “手工作坊” 模式，同时兼顾开发灵活性、运行性能与跨平台能力。

它的核心是多级分块（Tiling），将庞大计算任务切分成标准化的数据块。开发者只需声明 “每块数据在哪里计算”，编译器自动完成数据搬运、线程同步等底层工作，性能直逼手写 CUDA 的硬件上限。这本质上是把资深工程师手工调优的成熟思路抽象为标准化编程原语。TileLang 主攻大模型核心算子：通用矩阵乘法、反量化 GEMM、FlashAttention、线性注意力、稀疏 MLA。这些算子承载大模型训练和推理 90% 以上计算量，直接决定 AI 芯片算力利用率。算子速度每提升一倍，都意味着大模型训练成本下降、推理效率提升，从而大幅摊薄算力成本。DeepSeek 能够把 API 价格打到行业极致，除了模型架构本身的精妙设计，还因为底层拥有这套极高效的自研算子工具链 TileLang。

站在 TVM 肩膀上，性能跑进全球第一梯队
从技术角度而言，TileLang 没有从零去写一套完整编译器，它复用了Apache TVM现成的编译能力。Apache TVM 是华盛顿大学陈天奇团队开发的开源机器学习编译器框架，现为 Apache 顶级项目，在 AI 界的地位相当于 LLVM 在编程语言界的地位。它的作用是 AI 芯片与大模型之间的“万能翻译官兼超级优化师”，不管上层是什么框架，底层是什么芯片，它都能编译成适合该芯片的高性能机器码。TileLang 在 TVM 之上包装了一层简单好用的语法，专门处理大模型里最重要的分块计算。开发者不用关心底层硬件细节，只用很少代码，就能把芯片算力跑到接近上限。TileLang 最早由北京大学团队在 TVM 编译器基础设施上孵化。DeepSeek 把它接过去、推到生产，再反向捐赠给开源社区。截至 2026 年 9 月，仓库收获 7.6k stars、1,992 次 commit，已经成为事实上的国产 DSL 标杆。它的效果可以用官方基准数据说话：基于 TileLang 优化的 DeepSeek V3.2 稀疏注意力 TopK 算子，比原生 PyTorch 快 2–20 倍；其推理中最核心的 MLA、FlashAttention、LinearAttention 等算子，也已全部提供可直接商用的极致优化实现。同样在英伟达 H100 上运行核心算子，TileLang 的速度大幅超越了 Meta 的 PyTorch 原生实现。要知道，PyTorch 是目前全球使用人数最多、生态最庞大的 AI 核心框架，绝大多数主流大模型都基于它搭建。同时，它也击败了 OpenAI 的主力加速工具 Triton。Triton 是 OpenAI 倾力打造的、全球开源社区公认最主流的第三方加速标杆，世界顶级的大模型公司都在用它来压榨英伟达芯片的算力。

图注：H100 上各底层框架的 MLA 解码性能对比，相比传统和通用框架，FlashMLA 与 TileLang 具备极佳的硬件算力释放效率它甚至追平了英伟达官方调校的 cuBLAS 以及官方版 FlashAttention，直接进入全球顶尖性能的第一梯队。要注意的是，英伟达官方工具只支持自家芯片，而 TileLang 具有跨平台能力，在 AMD 顶级芯片 MI300X 上跑，同样拿到了接近硬件极限的跑分数据。

TileLang 架构深度拆解：为什么一套代码能跑遍所有芯片
TileLang 能用同一套编译器，灵活切换不同的芯片，核心在于从设计之初就采用了“后端解耦，目标与执行分离”的架构。传统编译器的痛点在于，把 “为哪款芯片生成指令” 和 “怎么编译加载运行” 两件事深度绑定。换一款芯片，整套逻辑都要推倒重来。而TileLang 把这两件事彻底拆开：

目标后端：只负责定义这款硬件的指令语法、优化规则，是硬件专属部分；
执行后：只负责通用的编译、加载、启动流程，和具体硬件无关。

每当新芯片适配时，执行后端完全不用改动，只需要新增一个对应的目标后端即可。每个目标后端都遵循统一的四层流水线结构，上层通用代码输入后，会逐步翻译成对应芯片的底层指令：第一层是语言方言（Dialect），把代码翻译成每家芯片能听懂的指令像针对英伟达 CUDA 后端：它负责翻译并注入 WGMMA 和 TCGEN05 等张量核心指令；针对 AMD ROCm 后端：它对应翻译为 MFMA 和 WMMA 核心架构指令；针对华为昇腾后端：它则负责翻译为 Ascend C 的底层向量指令与矩阵运算原语。第二层是上下文贡献（Context）。这一层要感知你当前跑的是什么芯片，并把不同硬件的规格参数统一转换成标准格式；记录整个编译过程的状态，保证前后一致。第三层是Pass 流水线（Pass Pipeline），这是整个流程最核心的翻译优化环节，要把高层逻辑转换成芯片能跑的底层代码，并且针对不同芯片做不同优化。第四层是主机 / 设备代码生成（Codegen）层，把优化好的代码拆成两份，主机侧代码运行在 CPU 上，负责任务调度；设备侧代码跑在 AI 加速器上，承担核心计算，两段代码组合完成最终编译与执行。截至目前，TileLang 覆盖主流 AI 算力平台：英伟达 CUDA 是核心主力，拥有最完整的功能和优化；AMD ROCm、苹果 Metal、华为昇腾 950 为官方支持级；LLVM CPU、WebGPU、CuTe DSL 等处于实验阶段；沐曦、摩尔线程、海光等国产芯片则通过生态项目适配，基本覆盖了国内外主流 AI 算力平台。这套解耦架构不仅解决了跨硬件适配的难题，还向上开放了精细化的硬件控制力。开发者只用接近 Python 的简洁语法，就能对底层硬件做精准操控，最终跑出媲美手写 CUDA 的性能。控制算子运行效率的核心，在于三个维度：存储控制、线程调度、并行节奏。TileLang 允许开发者直接在高级的 Python 代码中，对这三个底层维度进行显式控制。

存储控制（数据放哪）：芯片内部有三层存储。全局内存容量最大，但速度最慢；共享内存速度中等；寄存器容量最小，访问速度最快。TileLang 允许开发者在 Python 中直接用代码决定数据放在哪一层存储里，避免计算单元因为等待数据而闲置。
线程调度（任务怎么分）：TileLang 通过一句代码，就能把庞大的矩阵计算像 “划分网格” 一样，精准分配给芯片里成千上万个计算核心。例如把一个巨大的计算切分成 100 份任务，在芯片上拉起 100 组计算单元（线程块），每组里面跑 128 条并发工序（线程），多路并行同时开工。
并行节奏（传输和计算怎么并发）：TileLang 用一句 T.Pipelined(num_stages=3) 直接开启三级软件流水线，让数据搬运和算力计算异步并行，一边搬数据一边算数据，绝不让芯片闲着。

用这三个维度编写的通用矩阵乘法（GEMM）算子，可以把传统 CUDA 实现通常需要 500 行以上的代码，直接压缩到 30 多行 Python 代码。

押注华为昇腾，DeepSeek 最大的赌注
9 月 21 日，据《The Information》报道，在投资人闭门会上，梁文锋将“使用华为或其他国产芯片训练模型”描述为公司当前“最大的赌注之一”，并明确表示此役“必须成功”。据彭博社此前报道，DeepSeek 已投入 25.6 亿美元的订单，采购至少 16 万颗华为昇腾 950DT 加速器，用于部署在内蒙古乌兰察布在建的吉瓦级数据中心，这笔订单最快 2026 年 Q4 开始交付。这批昇腾 950DT 主要面向推理侧，训练环节目前仍主要依赖英伟达。在算力布局上，DeepSeek 没有将国产芯片视为英伟达的退路，而是被摆在了核心战略位。梁文锋的判断是，国产芯片走向全球一线只是时间问题，DeepSeek应该率先行动，适应英伟达之外的半导体硬件。DeepSeek 转向昇腾，其背后仍有三个难关。第一块：训练软件栈全部重写。DeepSeek-V3、V4 的训练代码，最早是为英伟达的 Tensor Core / Hopper 架构 WGMMA 指令写的。要在昇腾上跑，必须把每一处 cuBLAS / NCCL 调用替换成 Ascend C / HCCL。同时，FlashMLA、DeepEP、DeepGEMM 这些自研核心算子，要逐一在昇腾 950 上验证性能。第二块：通信原语底层替换。大模型训练要在成千上万颗加速器之间同步梯度，通信原语是并行训练的基石。英伟达生态的 NCCL 与昇腾体系的 HCCL，在 API 设计、性能曲线、容错模型上完全不同。从 NCCL 到 HCCL 的迁移不是改个 import 那么简单，任何一个集合通信操作的适配偏差，都可能导致整轮训练静默失效。第三块：芯片产能押注。更隐蔽的风险藏在供应链。推动 DeepSeek 转向国产芯片的动因是美国出口管制，而同一管制政策也制约着华为扩大 Ascend 950 产能的能力。DeepSeek 押注 16 万颗昇腾的隐含前提是：华为的 Q4 交付不能掉链子。目前，DeepSeek 正在训练一款 2 万亿参数的大模型，参数规模超过现有旗舰 V4 的 1.6 万亿；梁文锋同时透露，公司已规划 8 万亿参数的更大模型。模型规模越大，训练与推理对算力的需求就呈指数级增长。据梁文锋此前估算，训练一款与 OpenAI 同级别的大模型，大约需要 5 万颗英伟达 GB300 处理器，对应需要 20 万颗华为昇腾 950 芯片。尽管 DeepSeek 全力推进国产算力迁移，但现实是华为当前的芯片供货仍有缺口。知情人士透露，受先进内存等核心元器件短缺影响，华为正面临生产瓶颈，DeepSeek 短期内仍无法完全摆脱对英伟达硬件的依赖。为此，DeepSeek 已完成多轮大规模融资，持续投入算力扩张。未来 12 个月里，至少有 4 个变量会决定这场赌局的胜负：华为 Ascend 950 的 Q4 交付率，决定 DeepSeek 训练时间表；TileLang 昇腾后端是否能在 6 个月内追平 CUDA 后端，决定开发效率；字节、腾讯、阿里是否跟随押注昇腾，决定生态规模；美国出口管制是否再次升级，决定这条路的下限。这些问题里，还没有定论。但一个趋势已经清晰，TileLang 的开源，把“国产算力生态怎么建”这件事，从一个工程问题变成了一个生态问题。对比 CUDA 与 TileLang，两者代表了两种完全不同的生态路线：CUDA 是闭源、全栈、硬件绑定的厂商主导生态，积累深厚但迁移成本极高；TileLang 是开源、聚焦、跨硬件的社区化工具，专注解决大模型最核心的算子开发问题，灵活度高、移植成本低。TileLang 的意义不在于替代 CUDA，而是补上了当下最紧迫的短板，让国产芯片不需要从零搭建完整的软件生态，就能快速拥有大模型训练所需的顶级算子能力。当越来越多的大模型厂商基于 TileLang 开发算子，国产芯片的适配成本会指数级下降，整个生态的迭代速度也会大幅加快。某种意义上，DeepSeek证明了用国产芯片，一样能堆出世界级规模的超算集群。当大模型竞争进入算力成本与供应链安全的深水区，“用软件定义算力” 不再是一句口号，TileLang 这样的底层工具创新，正在悄然重构国产算力生态的格局。参考链接：https://github.com/tile-ai/tilelanghttps://www.theinformation.com/articles/deepseek-bets-big-huawei-chips-bypass-u-s-export-controls?utm_campaign=Editorial&utm_content=Article&utm_medium=organic_social&utm_source=threads,twitter&__cf_chl_rt_tk=iW7ZBFXRrEVuB9NBtCSW3X3dfxuUsmhAeiCUuxwa9RI-1790041475-1.0.1.1-NZ2.jEQlxsBBSsJDT6m6kKkiC9ImC8gHkVtqSlKM8iI