tile-ai/tilelang

GitHub Trending (daily) 工具

摘要

TileLang 是一个基于 TVM 构建的开源领域专用语言与编译器,用于编写高性能 GPU/CPU/NPU 内核,例如 GEMM 和 FlashAttention。近期新增了对 Huawei Ascend 950 的支持、LSP,以及面向 CUDA、ROCm 和 Metal 的多后端方言。

一个领域专用语言,旨在简化高性能 GPU/CPU/加速器内核的开发。
查看原文
查看缓存全文

缓存时间: 2026/10/01 15:36

tile-ai/tilelang

来源:https://github.com/tile-ai/tilelang

Tile Language

PyPI 版本 (https://badge.fury.io/py/tilelang) Ask DeepWiki (https://deepwiki.com/tile-ai/tilelang) Discord (https://discord.gg/TUrHyJnKPG) Puzzles (https://github.com/tile-ai/tilelang-puzzles) LSP (https://github.com/tile-ai/tilelang-lsp) 文档 (https://tilelang.com/) · 安装 (https://tilelang.com/get_started/Installation.html) · 示例 (https://github.com/tile-ai/tilelang/tree/main/examples) · 发布版本 (https://github.com/tile-ai/tilelang/releases) · 贡献指南 (https://github.com/tile-ai/tilelang/blob/main/CONTRIBUTING.md) · TileLang (Ascend 950)

Tile Language(tile-lang)是一种简洁的领域专用语言,旨在简化高性能 GPU/CPU/NPU 内核(如 GEMM、Dequant GEMM、FlashAttention、LinearAttention)的开发。通过采用 Pythonic 语法,并在 TVM (https://tvm.apache.org/) 之上构建底层编译器基础设施,tile-lang 让开发者在不牺牲前沿性能所需的底层优化的同时,专注于提升生产力。

最新动态

  • 2026-09-30 — Ascend 950 后端: TileLang 现已正式支持华为 Ascend 950 NPU,具备原生代码生成、自动调度与同步、SIMD/SIMT 向量编程等功能。欢迎探索 Ascend 的 GEMM、FlashAttention 等示例。
  • 2026-08-04 — TileLang LSP 开源 (https://github.com/tile-ai/tilelang-lsp): 发布了 TileLang 的 Language Server Protocol 实现,支持 buffer 形状、数据类型、作用域和推断布局的内联提示(inlay hints),以及悬停详情和精确诊断。
  • 2026-08-03 — TileLang v0.1.13 (https://github.com/tile-ai/tilelang/releases/tag/v0.1.13): 发布了多后端语言方言、编译器诊断中的源码位置信息、新的 CUDA 和 Metal 硬件路径,以及大量正确性修复。此版本移除了若干遗留 API;升级前请阅读兼容性说明。
  • 2026-07-30 — SM120 NVF4 block-scaled MMA (https://github.com/tile-ai/tilelang/pull/2364): 为 T.mma_gemm_blockscaled 添加了优化的 Blackwell 路径,并提供了相应的 SM120 示例。
  • 2026-07-28 — Metal 4 协作张量 GEMM (https://github.com/tile-ai/tilelang/pull/2252): 为 Apple M5 添加了协作张量 T.gemm 支持,同时保留了不支持的形状和系统的 simdgroup 回退方案。

更早的动态(2025–2026)

2026

  • 2026-07-28 — 源码感知的编译器诊断 (https://github.com/tile-ai/tilelang/pull/2751): 将 Python 源码位置信息带入 TIRX,并在编译器错误中展示。
  • 2026-07-24 — 多后端语言方言 (https://github.com/tile-ai/tilelang/pull/2734): 围绕共享语义重构了语言层,并提供了静态的 CUDA、ROCm 和 Metal 方言。
  • 2026-07-23 — 面向 dLLM 的 Block-causal 注意力 (https://github.com/tile-ai/tilelang/pull/2499): 为扩散语言模型添加了定长和变长 block-causal 注意力示例。
  • 2026-07-22 — IR Lower Trace (https://github.com/tile-ai/tilelang/pull/2725): 引入了一个调试工具,用于查看编译器每个 pass 以及最终代码生成步骤中的 IR 变化。
  • 2026-07-22 — DeepSeek V3.2 稀疏 MLA 反向传播 (https://github.com/tile-ai/tilelang/pull/2592): 根据 head-block 大小自适应选择 launch 宽度。
  • 2026-07-21 — DeepSeek V3.2 top-k 优化 (https://github.com/tile-ai/tilelang/pull/2659): 改进了稀疏注意力 top-k 选择器的内存访问模式,在报告的基准测试中性能提升约 1.9 倍。
  • 2026-07-16 — 编译器 pass 计时 (https://github.com/tile-ai/tilelang/pull/2622): 为编译器 pass 添加了性能剖析,并支持可配置的报告阈值。
  • 2026-07-12 — IKET profiler 集成 (https://github.com/tile-ai/tilelang/pull/2515): 添加了 CUDA timeline 插桩和性能剖析支持。
  • 2026-07-08 — TileLang v0.1.12 (https://github.com/tile-ai/tilelang/releases/tag/v0.1.12): 添加了 LLVM 后端、tile 调度器、后端注册表、pass 可视化器,并扩展了 Blackwell 支持。
  • 2026-07-06 — Pass Visualizer (https://github.com/tile-ai/tilelang/pull/2449): 引入了用于查看编译器变换的结构树浏览器。
  • 2026-06-26 — 跨主机 CUDA 二进制缓存 (https://github.com/tile-ai/tilelang/pull/2459): 使已编译的 CUDA 二进制文件可在兼容主机之间复用。
  • 2026-06-24 — Tile 调度器 (https://github.com/tile-ai/tilelang/pull/2441): 为内核作者引入了持久化 tile 调度原语。
  • 2026-06-24 — 后端 CodeGen 注册表 (https://github.com/tile-ai/tilelang/pull/2442): 将设备端和主机端 CodeGen (https://github.com/tile-ai/tilelang/pull/2446) 分发逻辑移至后端注册表之下。
  • 2026-06-18 — LLVM 后端 (https://github.com/tile-ai/tilelang/pull/2409): 通过 LLVM 添加了 CPU 降低(lowering)和执行支持。
  • 2026-06-18 — 任意布局 TMA 降低 (https://github.com/tile-ai/tilelang/pull/2380): 为 swizzled 共享内存布局启用了 TMA 传输。
  • 2026-06-16 — Pass Diff (https://github.com/tile-ai/tilelang/pull/2375): 添加了编译器 pass 之间的 IR 对比功能,用于调试 lowering 变化。
  • 2026-06-08 — TileLang v0.1.11 (https://github.com/tile-ai/tilelang/releases/tag/v0.1.11): 扩展了 scan、pipeline、后端、CUDA、ROCm 和 Metal 功能。
  • 2026-05-25 — Scan 算子 (https://github.com/tile-ai/tilelang/pull/2262): 引入了 tile 级别的 scan 原语。
  • 2026-05-25 — TileLang v0.1.10 (https://github.com/tile-ai/tilelang/releases/tag/v0.1.10): 扩展了 AMD 和 Blackwell 支持,添加了初步的 Metal GEMM,改进了 Windows 打包,并扩展了自动调优功能。
  • 2026-05-24 — CDNA4 MXFP4 (https://github.com/tile-ai/tilelang/pull/2132): 为 AMD gfx950 添加了 FP4 E2M1 矩阵核心支持。
  • 2026-05-22 — Metal simdgroup GEMM (https://github.com/tile-ai/tilelang/pull/1869): 添加了首个使用 simdgroup_matrix MMA 的 Metal T.gemm 路径。
  • 2026-05-20 — Cluster 拷贝 (https://github.com/tile-ai/tilelang/pull/1908): 引入了 T.copy_cluster,用于 TMA 多播和 SM 到 SM 的 cluster 传输。
  • 2026-05-20 — TMA gather/scatter (https://github.com/tile-ai/tilelang/pull/2129): 添加了 tile::gather4 和 tile::scatter4 支持。
  • 2026-05-20 — 原生 SM75 MMA GEMM (https://github.com/tile-ai/tilelang/pull/2198): 为 Turing GPU 添加了 FP16、INT8 和 INT4 tensor core 路径。
  • 2026-05-20 — TIRX 迁移 (https://github.com/tile-ai/tilelang/pull/2216): 将 TileLang 的 IR 用法迁移到 TVM 的 TIRX 表示。
  • 2026-05-11 — 并行自动调优 (https://github.com/tile-ai/tilelang/pull/2159): 添加了流水线编译、分组编译和多 GPU 基准测试功能。
  • 2026-05-07 — DeepSeek V4 算子 (https://github.com/tile-ai/tilelang/pull/2148): 为 DeepSeek V4 工作负载添加了 TileLang 示例。
  • 2026-05-06 — Windows 支持 (https://github.com/tile-ai/tilelang/pull/2093): 添加了完整的 Windows 构建和运行时支持,并进行了跨平台修复。
  • 2026-04-28 — MXFP8 分组 GEMM (https://github.com/tile-ai/tilelang/pull/2098): 在 Blackwell 上添加了支持转置 B 的 block-scaled 分组 GEMM 示例。
  • 2026-04-25 — HISA 稀疏注意力索引器 (https://github.com/tile-ai/tilelang/pull/2069): 添加了分层稀疏注意力索引示例。
  • 2026-04-24 — Blackwell MXFP8 block-scaled GEMM (https://github.com/tile-ai/tilelang/pull/1945): 在 SM100 上添加了 MXFP8 block-scaled 矩阵乘法。
  • 2026-04-22 — TileLang v0.1.9 (https://github.com/tile-ai/tilelang/releases/tag/v0.1.9): 发布了 CuTe DSL GEMM V2、Metal 代码生成改进,以及无需主机工具链的构建支持。
  • 2026-04-22 — RDNA3/RDNA3.5 WMMA (https://github.com/tile-ai/tilelang/pull/2044): 为 AMD gfx11 GPU 添加了 WMMA 降低支持。
  • 2026-04-20 — INT4 T.gemm (https://github.com/tile-ai/tilelang/pull/2063): 在 CUDA GEMM 路径中添加了 INT4 矩阵乘法。
  • 2026-04-17 — CUDA 源码内核 (https://github.com/tile-ai/tilelang/pull/1970): 引入了 T.CUDASourceCodeKernel,用于嵌入自定义 CUDA 源码。
  • 2026-04-15 — AutoDD 冻结区域 (https://github.com/tile-ai/tilelang/pull/2045): 添加了 __freeze__ 注解,用于在自动 delta debugging 期间保留选定的代码。
  • 2026-03-27 — TMA 存储 (https://github.com/tile-ai/tilelang/pull/1981): 为 T.tma_copy 添加了存储支持。
  • 2026-03-24 — 双 SM Blackwell 内核 (https://github.com/tile-ai/tilelang/pull/1882): 添加了双 SM TMA、TMEM 和 TCGEN5 MMA 支持。
  • 2026-03-23 — AMD RDNA4 (https://github.com/tile-ai/tilelang/pull/1951): 升级了 ROCm 路径并添加了 RDNA4 GPU 支持。
  • 2026-03-22 — SM100 上的 FlashAttention (https://github.com/tile-ai/tilelang/pull/1910): 添加了 Blackwell FlashAttention 示例。
  • 2026-03-18 — 生产者-消费者 warp 特化 (https://github.com/tile-ai/tilelang/pull/1909): 添加了自动 warp 特化流水线和 T.tma_copy API。
  • 2026-03-12 — Eager 模式自动调优 (https://github.com/tile-ai/tilelang/pull/1906): 为 eager JIT 内核启用了自动调优器。
  • 2026-03-10 — CPU T.gemm (https://github.com/tile-ai/tilelang/pull/1904): 为 CPU 目标添加了矩阵乘法支持。
  • 2026-03-05 — IR dump 配置 (https://github.com/tile-ai/tilelang/pull/1903): 添加了 TileLang pass 配置,用于转储中间 IR。
  • 2026-02-28 — CUDA cluster 原语 (https://github.com/tile-ai/tilelang/pull/1874): 添加了 cluster 启动、查询、同步和屏障操作。
  • 2026-02-28 — TCGEN5 MMA tensor-shared 路径 (https://github.com/tile-ai/tilelang/pull/1866): 添加了 tensor memory/shared memory 的 Blackwell GEMM 路径。
  • 2026-02-24 — 免主机工具链构建 (https://github.com/tile-ai/tilelang/pull/1833): 在支持预构建产物的情况下,无需主机 C/C++ 工具链即可安装。
  • 2026-02-23 — CuTe DSL GEMM V2 (https://github.com/tile-ai/tilelang/pull/1855): 为 CuTe DSL 后端添加了 SM90 和 SM100 GEMM V2 支持。
  • 2026-02-16 — TileLang v0.1.8 (https://github.com/tile-ai/tilelang/releases/tag/v0.1.8): 发布了动态流水线改进、日志文档、更丰富的布局表示,以及 AMD 修复。
  • 2026-02-14 — 跨 CUDA 版本的发布 wheel (https://github.com/tile-ai/tilelang/pull/1826): 将多个 CUDA 版本统一到单一 wheel 中。
  • 2026-02-14 — 分层归约 (https://github.com/tile-ai/tilelang/pull/1762): 添加了分层和 warp 级别的归约内建函数。
  • 2026-02-09 — CUDA runtime stubs (https://github.com/tile-ai/tilelang/pull/1821): 为 CUDA 11、12 和 13 兼容的 wheel 添加了延迟加载的 CUDART 和 NVRTC stubs。
  • 2026-02-08 — 布局可视化改进 (https://github.com/tile-ai/tilelang/pull/1811): 改进了 TileLang 布局的渲染和查看体验。
  • 2026-02-02 — TileLang Puzzles (https://github.com/tile-ai/tilelang-puzzles): 发布了十个难度递增的练习题,用于交互式学习 TileLang。

2025

  • 2025-12-18 — CuTe DSL 后端 (https://github.com/tile-ai/tilelang/pull/1421): 添加了通过 NVIDIA CUTLASS CuTe DSL 进行编译的支持;请关注 issue #1454 (https://github.com/tile-ai/tilelang/issues/1454) 中的后续进展。
  • 2025-12-17 — Z3 集成 (https://github.com/tile-ai/tilelang/pull/1367): 将基于 SMT 的符号推理集成到 TVM 的算术分析器中。
  • 2025-10-31 — Apache TVM FFI 迁移 (https://github.com/tile-ai/tilelang/pull/1108): 将运行时接口迁移至 apache-tvm-ffi,以降低主机端开销。
  • 2025-10-30 — TileLang v0.1.6.post2 (https://github.com/tile-ai/tilelang/releases/tag/v0.1.6.post2): 发布了最后一个兼容 Python 3.8 的 TileLang 版本。
  • 2025-10-07 — Apple Metal 后端 (https://github.com/tile-ai/tilelang/pull/799): 为 Apple Silicon 引入了 Metal 设备支持。
  • 2025-09-29 — Huawei Ascend 适配器 (https://github.com/tile-ai/tilelang-ascend): 在外部 TileLang Ascend 项目中发布了 AscendC 和 Ascend NPU IR 后端工作。
  • 2025-07-04 — 2:4 稀疏 tensor core (https://github.com/tile-ai/tilelang/pull/526): 引入了 T.gemm_sp 用于结构化稀疏矩阵乘法。
  • 2025-06-05 — NVRTC 执行后端 (https://github.com/tile-ai/tilelang/pull/461): 添加了 NVRTC 路径,以减少生成 CUDA 模板的编译时间。
  • 2025-04-14 — AMD MI300X 上的 FlashMLA (https://github.com/tile-ai/tilelang/tree/main/examples/deepseek_mla/amd): 发布了优化的 AMD 实现及相关文档。
  • 2025-03-03 — H100 上的 MLA 解码 (https://github.com/tile-ai/tilelang/tree/main/examples/deepseek_mla): 发布了简洁的 TileLang 实现、基准测试和优化过程讲解。
  • 2025-02-15 — WebGPU 代码生成 (https://github.com/tile-ai/tilelang/pull/86): 添加了初步的 WebGPU 后端。
  • 2025-02-12 — TileLang v0.1.0 (https://github.com/tile-ai/tilelang/releases/tag/v0.1.0): 发布了首个 v0.1 公开版本。
  • 2025-02-10 — 调试和布局工具 (https://tilelang.com/tutorials/debug_tools_for_tilelang.html): 添加了 T.print 和 fragment 布局可视化工作流。
  • 2025-01-20 — TileLang 开源 (https://github.com/tile-ai/tilelang): 项目正式向公众开放。

完整的变更日志和兼容性说明请参见所有发布版本 (https://github.com/tile-ai/tilelang/releases)。

平台与后端支持

TileLang 正在演进为一个多后端编译器(TileLang-X),围绕模块化的后端抽象构建。关于设计详情,请参阅后端架构;当将 TileLang 移植到新的后端时,请让编码 agent 使用后端集成技能。

当前支持的后端如下所列。Primary 表示 TileLang 的核心后端,Supported 和 Experimental 后端在主仓库中实现。Ecosystem 适配器位于独立仓库中,不包含在 TileLang 发布 wheel 中,且可能遵循独立的兼容性计划。

预构建 wheel 适用于 Linux x86-64/AArch64、Windows x86-64 和 macOS arm64。

TileLang 使用 Target 对象表示编译目标。默认的 auto 目标可检测 CUDA、HIP、Metal 和 Ascend 设备;当为其他后端或架构编译时,请选择显式目标。

关于目标语法、架构选项和后端特定说明,请参阅目标指南 (https://tilelang.com/get_started/targets.html);关于安装和已测试设备的详细信息,请参阅相应的适配器仓库。

后端Target平台和硬件支持级别备注
NVIDIA CUDAcudaLinux x86-64/AArch64、Windows x86-64;从 SM70 到 SM120 的代码路径Primary包含在发布 wheel 中并有 CI 覆盖;TMA、WGMMA 和 TMEM 功能需要相应的 GPU 架构。
AMD ROCm/HIPhipLinux;CDNA 和 RDNA GPU,包括 gfx942/gfx950 路径Supported包含在 Linux wheel 中;需要 ROCm 运行时。CI 在自托管的 gfx942 (MI300X) runner 上运行;gfx950 尚未覆盖。
Huawei Ascend 950ascendLinux;Ascend 950 NPUSupported使用 USE_ASCEND=ON 从源码构建;需要 CANN 和 torch_npu。请参阅 Ascend 指南。
Apple MetalmetalApple Silicon 上的 macOSSupported包含在发布 wheel 中并有 CI 覆盖;Metal 4 协作张量在受支持的 M5 系统上可用。
LLVM CPUllvm主机 CPUExperimental使用 USE_LLVM=ON 从源码构建;需要 LLVM 15 或更高版本。
NVIDIA CuTe DSLcutedslNVIDIA GPUExperimental需要…

相似文章

@seclink: 有点意思 ....

X AI KOLs Timeline

TileRT 是一款基于 tile 的运行时,能够实现超低延迟的 LLM 推理,近期里程碑包括在万亿参数模型上达到每秒 1000+ tokens。它支持 DeepSeek-V3.2 和 GLM-5 等模型,并在 GitHub 上开源提供。