---
url: /tilelang/introduction.md
description: >-
  TileLang 是什么？这门让 DeepSeek V4 系列大部分算子落地、2026 年 9 月 30 日正式支持华为昇腾 950 的算子 DSL，最新
  v0.1.15，一条 pip install tilelang 就能写 GEMM 和 FlashAttention。一文讲清 TileLang
  的定位、多后端架构、今天 DeepSeek 开源的昇腾六件套、真实性能数据与上手姿势
---

# 认识 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](https://github.com/tile-ai/tilelang) |
| 官网 | [tilelang.com](https://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 源码
```

### 下一步去哪儿

1. 先把官方 [examples](https://github.com/tile-ai/tilelang/tree/main/examples) 的 elementwise 和 GEMM 跑一遍
2. 去做 [TileLang Puzzles](https://github.com/tile-ai/tilelang-puzzles) 那 10 道题，做完基本就算入门了
3. 装 [TileLang LSP](https://github.com/tile-ai/tilelang-lsp)，没它你会写得很痛苦
4. 真要上生产，研究 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 这类东西的普及程度，才是国产算力真正能不能用起来的温度计。**

## 进阶

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