认识 TileLang —— DeepSeek 拿来写算子的那门语言，今天把昇腾也拿下了

先说一个让我有点难受的事实。

这两年芯片的算力一路狂飙，一张卡从几百 TFLOPS 干到上千 TFLOPS，集群从千卡卷到十万卡。但能把这些算力真正榨干的人，数量几乎没变——全球能手写高性能融合算子的工程师，撑死几千号人。剩下的绝大多数人，包括我，写的 PyTorch 代码跑在硬件上大概只吃了三成性能，剩下七成在显存带宽上空转。

更难的是国产芯片。昇腾有算力，但你要用就得写 Ascend C，那玩意儿的学习曲线不是陡，是垂直的。于是就出现了一个荒诞的局面：卡买回来了，跑不满；不是卡不行，是没人会写。

然后今天，2026 年 9 月 30 日，DeepSeek 一口气开源了面向昇腾的六件套：

TileLang 昇腾版（昇腾 950 后端，官方支持，原生代码生成）
DeepGEMM-Ascend（矩阵乘，跑到硬件极限的 99.8%）
DeepEP-Ascend（MoE 专家并行通信，EP8 dispatch 375 GB/s）
TileKernels v2.0.0（一套 Python 接口，同时跑 NVIDIA 和昇腾）
FlashMLA（稀疏注意力 / MLA 算子）
DeepSelect（TopK 算子，比 torch.topk 快 2～20 倍）

六个仓库，全部 MIT（TileLang 本体除外，见后文）。而且 DeepSeek 的说明里有一句话特别扎心："目前其训练中用到的每一个 TileLang 算子，在昇腾上都有对应的高性能实现。"

翻译一下：不是"我们适配了一下"，是"我们训练 V4 用的算子，昇腾上一个不落，全部对齐了"。

芯片的战争打到今天，胜负手已经不在晶体管上了，在有没有一门让人写得出算子的语言。CUDA 赢不是因为英伟达的硬件设计得多优雅，是因为它先让人写得起。

而这篇的主角，就是那门语言：TileLang。

什么是 TileLang

一句话：TileLang（tile-lang）是一门用来写高性能 GPU / CPU / NPU 算子的领域专用语言（DSL）。

再具体一点：它长得很像 Python，底层编译器架在 TVM 上，你用 @tilelang.jit 装饰一个函数，在里头描述"数据怎么切块、怎么搬、怎么算"，它给你吐出能跑满硬件的 kernel。官方对自己的定位是——让你专注生产力，同时不牺牲那些为了达到 SOTA 性能所必需的底层优化。

注意这后半句，它是 TileLang 和很多"简化版"方案的根本区别。很多 DSL 的思路是"我帮你优化，性能差一点你认了"；TileLang 的思路是"布局（layout）、流水线（pipeline）、并行策略这些脏活我自己管，但你要下场手搓我也给你留口子"。

几个基本盘：

项目	情况
仓库	github.com/tile-ai/tilelang
官网	tilelang.com
最新版本	v0.1.15（PyPI，2026-09-30 发布）
Star / Fork	约 7810 / 781
首次开源	2025-01-20（v0.1.0 于 2025-02-12 发布）
安装	`pip install tilelang`
Python	3.10 ～ 3.14
底层	TVM（TIRX IR）

顺手纠个偏：TileLang 不是某个公司的内部工具外泄，它是独立的开源项目（tile-ai 组织），2025 年 1 月就公开了。DeepSeek 是它最重度的用户和贡献者之一，不是它的所有者。这两者的关系，别搞反。

核心特点

一、一套代码，多种芯片

这是今天最大的新闻，也是 TileLang 现在最硬的卖点。它正在往 TileLang-X 演进——一个围绕模块化后端抽象搭起来的多后端编译器。目前的后端全家福：

后端	Target	硬件	支持级别
NVIDIA CUDA	`cuda`	SM70 ～ SM120（Volta 到 Blackwell）	Primary
AMD ROCm/HIP	`hip`	gfx942 / gfx950（MI300X 等）	Supported
华为昇腾 950	`ascend`	Ascend 950 NPU	Supported（2026-09-30 新增）
Apple Metal	`metal`	Apple Silicon，M5 上支持 Metal 4 协同张量	Supported
LLVM CPU	`llvm`	主机 CPU	Experimental
NVIDIA CuTe DSL	`cutedsl`	NVIDIA GPU	Experimental
WebGPU	`webgpu`	WebGPU 运行时	Experimental

还有一批生态适配器（独立仓库，不在官方 wheel 里）：昇腾 A2/A3、沐曦 MetaX MACA、摩尔线程 MUSA、海光 HYGON、天数智芯 Sunrise TANG。

国内做芯片的厂商有多少家，这张表就还能再长多长。某种意义上，TileLang 正在变成国产算力的"通用插座"——芯片厂不用各自从零造一套编程栈，开发者也不用每换一款卡就重学一遍。

Target 默认写 auto 就行，它会自己探测 CUDA / HIP / Metal / Ascend。想给别的架构编译，显式指定即可。

二、Pythonic 语法，但管得住底层

TileLang 的核心抽象是**"块（tile）"**。你不用手写线程索引和共享内存地址，而是声明"我要一块 128×32 的 shared"、"一块 128×128 的 fragment"，剩下的搬运动作由编译器生成。

它替你管的脏活包括：

内存作用域：T.alloc_shared（片上共享内存）、T.alloc_fragment（寄存器），你只说想要什么，不操心地址
布局推断：共享内存的 swizzle、寄存器 fragment 的排布，编译器自己算，也提供可视化工具给你看
自动流水线和同步：T.Pipelined 一个 num_stages=3 就把"预取—计算—写回"叠起来；昇腾后端还带自动调度与自动同步
warp 专业化：生产者—消费者 warp specialization 自动做（2026-03-18 起）
自动调优：autotuner 会帮你把 block 大小、stage 数、线程数全搜一遍，比人肉试快得多

你想下场手搓？也留了口子：T.copy / T.gemm / T.reduce / T.tma_copy / T.copy_cluster 这些原语都在，T.CUDASourceCodeKernel 甚至允许你直接嵌一段自己的 CUDA 源码进去。

三、算子覆盖面是真的全

官方 example 列表基本就是大模型kernel的目录：

GEMM 与量化：GEMM、分组 GEMM、FP8 GEMM、反量化 GEMM、块缩放 GEMM
注意力与序列模型：FlashAttention、Flash Decoding、块稀疏注意力、线性注意力、GDN
模型级算子：DeepSeek MLA、DeepSeek V3.2、DeepSeek V4、DeepSeek mHC
架构特化：AMD、TCGEN05、SM120（Blackwell NVF4）

看到那行 "DeepSeek V4" 了吗？DeepSeek 训练 V4 系列时大部分算子是 TileLang 写的，这不是营销话术，是 GitHub 上白纸黑字的 News 条目（2026-05-07 "DeepSeek V4 operators"）加上今天的官方声明。

四、调试工具链意外地齐

写算子最痛苦的不是写，是写错了不知道错在哪。TileLang 在这块堆了不少东西：

TileLang LSP（2026-08-04 开源）：inlay hints 直接显示 buffer 的 shape、dtype、scope 和推断出的 layout，hover 看细节，诊断带源码位置
IR Lower Trace：看每一趟 compiler pass 之后 IR 变成什么样
Pass Visualizer / Pass Diff：结构树浏览 + IR 差异对比
Layout 可视化：把 fragment 布局画出来给你看
TileLang Puzzles：10 道由浅入深的练习题，交互式学语言（这个我吹爆，比干读文档强太多）
IKET profiler 集成：CUDA timeline 打点

今天的六件套，各自管什么

热点归热点，还是得拆开看清楚，不然容易把"DeepSeek 开源了"当成一件事。

组件	干什么	关键数据
TileLang（昇腾后端）	昇腾 950 原生代码生成 + 自动调度同步 + SIMD/SIMT 向量编程	2026-09-30 官方支持
DeepGEMM-Ascend	矩阵乘，API 完全兼容 NVIDIA 版 DeepGEMM	BF16 达硬件极限 99.8%（431 / 432 TFLOPS），FP4 98.3%（1701 / 1730）
DeepEP-Ascend	MoE 专家并行 all-to-all 通信	EP8 dispatch 373–375 GB/s、combine 345–347 GB/s，EP≤32 时达物理有效带宽的 90–95%
TileKernels v2.0.0	数十个算子：MoE 路由、量化、Engram、mHC、RoPE	同一套 Python API，NVIDIA 与昇腾运行时自动切后端，已用于 DeepSeek 内部训练与推理
FlashMLA	稀疏注意力 / MLA 算子，提升长上下文效率	支持 NVIDIA 与昇腾
DeepSelect	DSA 稀疏注意力与采样器的 TopK 算子	比 `torch.topk` 快 2～20 倍（`topk` ≤ 4096）

几个我觉得比数字更值得看的点：

第一，DeepGEMM-Ascend 那份性能表是真的狠。 4096×7168×16384 的 BF16 GEMM，实测 431 TFLOPS，硬件极限 432 TFLOPS——99.8%。FP8 也到 99.5%。这个数字的含义是：在这颗芯片上，你几乎不可能写得比它更好了。 顺带一提，它的 mHC kernel 就是 TileLang 写的，DeepGEMM-Ascend 自己的致谢里专门点了 TileLang。

第二，DeepEP-Ascend 的带宽数据要打个时间折。 官方明说了：这组数字是用华为提供给 DeepSeek 的 PoC HDK（概念验证版固件包）+ 手动配置跑出来的，不是公开发布版本。真正的商用版本（Atlas 850E 的 Q3 HDK）预计 2026 年 10 月 15 日左右开放。用旧 PoC 版本的人会看到更低的带宽。这种"把家丑写在 README 里"的做法，我respect。

第三，华为这边也在同步开源。 昇腾 950 之外，华为提供了联合定义的 SuperPoD Flex 超节点 + UBL128 组网：128 卡 3.2 Tbps 单层交换 Scale-up，两层交换 Scale-out 到 256K 卡，配 ASC-COMM 自定义通信编程库。相关实践（大 EP 低时延推理、单卡/单机部署、大规模训练、长上下文 KV Cache 池化、Agentic RL）已在 CANN 社区开源。华为还披露了 DeepSeek-V4.1-Flash 的部分离线推理数据：EP32、128K 上下文下，单卡 2469 output tokens/s（TPOT 5ms），10ms 时 5102 tokens/s（不含 serving 调度与框架负载均衡）。

这一整套东西的意义，不在于"昇腾追上了 H100"。而在于它把"在国产卡上训出前沿模型"从玄学变成了工程——路径有人趟过了，代码摆在那儿了，剩下的是复现成本，不是探索成本。

什么时候该用它

别一上来就"我要学算子开发"。TileLang 的适用场景其实很窄，但窄的地方它无可替代：

你在做模型训练/推理的性能优化，PyTorch 原生算子是瓶颈，且现成库（cuBLAS、FlashAttention）覆盖不到你的融合需求
你在国产卡 / 非 NVIDIA 硬件上落地，没有 CUDA 生态可用，又不想写 Ascend C
你是芯片厂商或编译器工程师，需要给自家硬件快速搭一套可用的软件栈（看看那张生态适配器表，沐曦、摩尔线程、海光、天数都在这条路上）
你在做研究，想快速验证一个新的注意力变体或量化方案能不能跑得快

不适合的情况也直说：如果你只是想跑个模型，vLLM / SGLang 已经够了，别来受这个罪。 算子开发的痛苦程度，和你离硬件的距离成正比。

初体验：十分钟写个能跑的 GEMM

环境要求很低：Python 3.10+，一块 NVIDIA 显卡（SM70 起步）就行。昇腾需要 Ascend 950 + CANN 9.2.0+ 并从源码构建（USE_ASCEND=ON），门槛高一些，先拿 CUDA 试水。

装

bash

pip install tilelang
python -c "import tilelang; print(tilelang.__version__)"   # 应输出 0.1.15

想尝鲜可以用 nightly：

bash

pip install tilelang --find-links https://tile-ai.github.io/whl/nightly

写

一个带 FP32 累加的 FP16 矩阵乘，核心逻辑不到 20 行：

python

import torch
import tilelang
import tilelang.language as T

@tilelang.jit
def matmul(M, N, K, block_M=128, block_N=128, block_K=32,
           dtype=T.float16, accum_dtype=T.float32):
    @T.prim_func
    def main(
        A: T.Tensor((M, K), dtype),
        B: T.Tensor((K, N), dtype),
        C: T.Tensor((M, N), dtype),
    ):
        # 按 (N 方向, M 方向) 网格切块，每块 128 线程
        with T.Kernel(T.ceildiv(N, block_N), T.ceildiv(M, block_M), threads=128) as (bx, by):
            A_shared = T.alloc_shared((block_M, block_K), dtype)   # 片上共享内存
            B_shared = T.alloc_shared((block_K, block_N), dtype)
            C_local  = T.alloc_fragment((block_M, block_N), accum_dtype)  # 寄存器

            T.clear(C_local)
            # num_stages=3：三级流水线，预取/计算/写回自动重叠
            for k in T.Pipelined(T.ceildiv(K, block_K), num_stages=3):
                T.copy(A[by * block_M, k * block_K], A_shared)
                T.copy(B[k * block_K, bx * block_N], B_shared)
                T.gemm(A_shared, B_shared, C_local)
            T.copy(C_local, C[by * block_M, bx * block_N])

    return main

kernel = matmul(1024, 1024, 1024)

a = torch.randn(1024, 1024, device="cuda", dtype=torch.float16)
b = torch.randn(1024, 1024, device="cuda", dtype=torch.float16)
c = kernel(a, b)

torch.testing.assert_close(c, a @ b, rtol=1e-2, atol=1e-2)
print("对上了")

你会注意到几件事：没有任何线程索引，没有 __syncthreads()，没有手写共享内存地址。T.Pipelined(..., num_stages=3) 一行就把软件流水叠好了。这就是"生产力"那三个字的兑现方式。

想看编译器到底干了什么：

python

print(kernel.get_kernel_source())     # 生成的 CUDA 源码

下一步去哪儿

先把官方 examples 的 elementwise 和 GEMM 跑一遍
去做 TileLang Puzzles 那 10 道题，做完基本就算入门了
装 TileLang LSP，没它你会写得很痛苦
真要上生产，研究 autotuner——手工调 block 大小和 stage 数，性价比极低

几句必须说的丑话

版本在快速迭代。 v0.1.13 明确"移除了若干 legacy API"，升级前务必读兼容性说明。0.1.x 这个版本号本身就说明了一切。
昇腾后端今天才落地。 官方标注是 Supported 但需源码构建，要有 CANN + torch_npu；其它昇腾世代（A2/A3）走的是另一个生态仓库，别混。
生态适配器不在 wheel 里。 沐曦、摩尔线程这些要单独装，且各自维护兼容节奏。
DeepEP / TileKernels / DeepGEMM 的部分能力仍在开发中。 DeepEP-Ascend 的 PP、Engram、Bucket 被标为实验性，reduce-scatter / all-reduce kernel 还没做完，专家负载均衡的通信 kernel 也没实现。TileKernels 只有 Ascend 实现支持 bf16。下单之前先看 README 里的 "Ongoing"。

最后一句个人立场：过去两年大家都在卷模型、卷参数、卷 Agent，很少有人愿意聊"算子"这种苦活累活。但模型再聪明，跑不动就是跑不动。TileLang 这类东西的普及程度，才是国产算力真正能不能用起来的温度计。

进阶

更多开源技术干货和学习资料，关注公众号「遇码」，领取专属福利。