← 所有项目

Contribution Tasks

tile-ai/tilelang

候选人的 GPU kernel / 训练性能专长与当前仓库的多后端扩展、transform pass 稳定性、JIT 缓存可观测性、autotune 等方向高度重合;但硬约束是没有任何 GPU,因此只能承担 CPU / Apple Silicon 上可验证的任务(LLVM CPU 后端编译与执行、Metal 后端编译产物验证、transform pass / IR dump / host_checks 等编译期路径、autotune 流程在 CPU 上的复现、文档与测试基建)。所有推荐任务都落在这些路径上。

当前方向:近 30 天提交热点集中在 testing/python(20)、tilelang/jit(6)、tilelang/contrib(5)、src/cuda(3)、tilelang/autotuner(2)、tilelang/ascend(1)、tilelang/language(1)、src/transform(1)、src/rocm(1)、src/tl_templates(1)。README 2026 年新闻主线是 LLVM CPU 后端(2026-06-18)、Metal simdgroup / cooperative-tensor GEMM(2026-05-22 / 2026-07-28)、Ascend 950 官方后端(2026-09-30)、parallel autotuning(2026-05-11)、IR Lower Trace / Pass Timing / Pass Diff / Pass Visualizer(2026-07 系列)、cross-host CUDA binary cache(2026-06-26)、v0.1.13→v0.1.15 每月一版。可安全承接的方向:CPU / Metal 后端 correctness 测试、transform pass 调试工具文档与示例、autotune 结构化 trial 报告的 CPU 复现、host_checks 静态检查扩展、JIT 缓存可观测性。

★ 8577Fork 8755 个候选任务Gemini:LongCat-2.0

更新于 2026-11-05T07:57:30+00:00 · 本次生成失败,展示上次结果 · 打开仓库 ↗

一、项目定位

Tile Language(tile-lang)是由 Tile-AI 维护的 Pythonic DSL,目标是让高性能 GPU/CPU/NPU kernel 的开发更接近写 PyTorch,同时仍能做 tile 级别的底层优化。它构建在 Apache TVM 之上,自 TVM 的 TIR 之上发展出 TIRX IR 与一套独立的 pass 管道,并在 IR lowering 之后对接多种后端:CUDA(含 CuTe DSL、NVRTC、SM75/90/100/120 各代 tensor core / block-scaled MMA / TMA / cluster copy / warp specialization)、ROCm/HIP(CDNA4、RDNA3/RDNA4、WMMA)、Metal(simdgroup、cooperative tensor GEMM)、Ascend 950 NPU(PTODSL、SIMT/SIMD 向量编程)、LLVM CPU、以及实验性的 WebGPU。项目当前处于活跃演进期,从 2025 年的 CuTe DSL 后端、Metal 起步,到 2026 年已覆盖 Blackwell(SM100/SM120)FP8/MXFP8/INT4、Windows 包、并行 autotune、LSP、源码级诊断、pass 可视化、Ascend 官方后端等。README 与 pyproject 自称 Beta(Development Status :: 4 - Beta),但多个新后端(WebGPU、部分 Ascend 能力、eager-mode autotune)仍在快速迭代。仓库规模可观:tilelang/ 下约 414 个 Python 文件、src/ 下约 398 个 C++ 文件、testing/ 约 499 个测试文件、examples/ 约 387 个示例、benchmark/ 与 maint/ 若干工具脚本。近 30 天内 release 节奏约每月一版(v0.1.13→v0.1.14→v0.1.15),提交热点集中在 testing/python、jit、contrib、autotuner、CUDA/Ascend/ROCm 后端与 transform pass,说明项目正处在多后端扩展与稳定性并重的阶段。

解决什么问题、给谁用

解决的核心矛盾是:手写高性能 kernel(GEMM、dequant GEMM、FlashAttention、MoE、稀疏 MLA 等)需要精细控制 tile、shared memory、tensor core / MMA 指令、TMA、pipeline,而传统 CUDA/Triton/CuTe 的开发与调试成本仍然很高。TileLang 让开发者在 Python 中用 T.copy、T.gemm、T.mma_*、T.copy_cluster、T.tma_copy 等 tile 级原语描述计算,由编译器自动完成 layout 推断、software pipeline、warp specialization、host 代码生成,并在多后端间复用同一套 kernel 描述。目标用户是:需要为特定硬件(H100/B200/Mi300/Apple Silicon/Ascend)自研高性能算子的性能工程师、训练框架底层开发者、编译器/DSL 研究者,以及希望快速验证新硬件特性的工程师。典型场景包括:为 DeepSeek V3/V4 / MLA / NSA 写自定义 attention kernel;为 MoE 写 fused MoE + quant/dequant 路径;为 Blackwell 写 block-scaled GEMM / TMA / two-SM 核;在 Apple Silicon 上用 Metal 验证算法;在 CPU/LLVM 上做 kernel 正确性回归;用 autotuner 自动搜优 tile/ pipeline 参数。

1. GPU 性能工程师:用 TileLang 替代手写 CUDA 或 Triton,开发 GEMM、attention、MoE、sparse attention 等算子,尤其适合需要跨代硬件(SM75→SM120)统一代码库的团队。2. 训练框架底层开发者:为自研 LLM 训练框架集成高性能 attention / fused MoE / block-scaled GEMM,利用 TileLang 的 autotune 和 pass 调试工具快速迭代。3. 编译器/DSL 研究者:研究 tile-level DSL、IR lowering、layout inference、software pipeline,TileLang 提供相对完整的端到端基础设施(Python 前端 → TIRX → 多后端 CodeGen)。4. 端侧 / 多硬件开发者:在 Apple Silicon 上通过 Metal 后端验证 kernel,或在 Ascend NPU 上利用 PTODSL 后端做算子开发,避免为每个硬件维护一套独立代码。5. 库维护者与贡献者:通过 maint/ 下的 audit 脚本、docs/ 的 CONTRIBUTING.md、format.sh、pre-commit 等工具参与。

同类项目与差别

核心能力

能力在哪成熟度
Pythonic tile-level DSL(T.copy / T.gemm / T.mma_* / T.copy_cluster / T.tma_copy / pipeline)tilelang/language/, tilelang/tileop/成熟
多后端 lowering(CUDA / ROCm / Metal / Ascend / LLVM CPU / WebGPU)src/backend/, src/cuda/, src/rocm/, src/metal/, src/ascend/, src/cpu/, src/webgpu/, tilelang/backend/成熟(CUDA/ROCm/Metal/Ascend/CPU 成熟;WebGPU 实验)
IR 基础设施(TIRX、pass 管道、IR Lower Trace、Pass Visualizer、layout 可视化)src/transform/, src/ir.cc, src/ir.h, tilelang/transform/, docs/compiler_internals/成熟
Autotuner(并行 autotune、eager-mode autotune、structured trial reporting)tilelang/autotuner/, tilelang/jit/成熟(eager-mode 实验)
Blackwell(SM100/SM120)block-scaled MMA、TMA、two-SM、TCGEN05examples/flash_attention_sm100, examples/blockscaled_gemm_sm100, examples/gemm_sm120, examples/gemm_tcgen05, examples/two_sm*, src/cuda/成熟
AMD CDNA4 / RDNA3/RDNA4(FP4、WMMA、ROCm 路径)examples/gemm_fp4, examples/amd/, src/rocm/, 3rdparty/hip-headers/成熟
Metal(simdgroup GEMM、cooperative tensor GEMM、Apple Silicon)tilelang/metal/, src/metal/, examples/gemm_metal, benchmark/matmul_metal成熟
Ascend 950 NPU(PTODSL、SIMT/SIMD、SIMT allreduce)tilelang/ascend/, src/ascend/, examples/ascend/, testing/ascend/成熟(官方支持,但仍在扩展)
LLVM CPU 后端(CPU T.gemm、host 代码生成)src/cpu/, tilelang/cpu/, examples/ 中的 CPU 示例成熟
LSP(Language Server Protocol、inlay hints、diagnostics)独立仓库 tile-ai/tilelang-lsp(README 与徽章引用)成熟
调试与诊断工具(源码级诊断、pass diff、layout 可视化、IKET profiler)tilelang/tools/, docs/tools/, maint/成熟
WebGPU 后端src/webgpu/, tilelang/webgpu/实验

阶段:项目处于快速成长期(Beta,v0.1.x),核心 CUDA 后端(SM75/90/100/120、TMA、warp specialization、CuTe DSL)相对成熟,Metal、ROCm、Ascend、LLVM CPU 后端在 2026 年密集补齐,WebGPU 与部分新特性(eager-mode autotune、block-scaled GEMM、scan operators)仍在实验/扩展期。从提交热点看,testing/python 占比最高,说明团队在积极补测试覆盖;jit、contrib、autotuner、多后端 codegen 的频繁提交表明功能仍在快速迭代。

技术栈:Python 3.10+(前端 DSL、autotuner、tools);C++(IR、transform、backend codegen,CMake >= 3.26);Apache TVM(TIRX、CodeGen、runtime,通过 3rdparty/tvm 子模块);PyTorch(torch、torch-c-dlpack-ext,作为数组后端与扩展构建依赖);z3-solver(layout 推断 / 符号求解);Cython + scikit-build-core(Python 扩展构建);CUDA toolkit(pip 版 nvidia-cuda-nvcc / nvrtc / curand,或 host CUDA);CUTLASS(3rdparty/cutlass,作为 CuTe DSL 后端);HIP headers(3rdparty/hip-headers,ROCm 路径);Ruff + clang-format + pre-commit(CI lint);GitHub Actions(.github/workflows);uv(推荐开发环境)

规模:Release:v0.1.15(2026-09-30)、v0.1.14(2026-09-02)、v0.1.13(2026-08-03),约每月一版。目录规模:tilelang/ 约 414 个 Python 文件、src/ 约 398 个 C++ 文件、testing/ 约 499 个测试文件、examples/ 约 387 个示例、benchmark/ 与 maint/ 若干工具。最近提交(2026-10-10–12)集中在 CUDA SM120 block-scaled MMA、CuTe DSL / NVRTC wrapper 修复、Ascend PTODSL allreduce、Autotune structured reporting、testing infrastructure 改进,显示多后端与稳定性仍在快速迭代。Docker 镜像覆盖 CUDA 11.8–12.8 与 ROCm。Star 数未在证据中提供,需验证。

二、架构与代码地图

TileLang 采用经典的编译器分层架构,自上而下分为五层。最上层是 接入层,以 tilelang.__init__.py 暴露的 tilelang.compile()、tilelang.jit 装饰器以及 tilelang.dsl / tilelang.language 中的 T.Kernel、T.Tensor、T.Pipelined、T.copy、T.gemm 等 Pythonic DSL 构成,用户在这一层写出看起来像 PyTorch 的 kernel 描述。DSL 原语会被捕获成 TIRX IR(TVM Script 风格的中间表示),缓存在 tilelang/cache 中,并由 tilelang/jit 负责即时编译与缓存命中。 往下一层是 调度层 / 编译器层,核心在 tilelang/transform、tilelang/analysis、tilelang/carver、tilelang/engine 以及 C++ 侧的 src/transform、src/ir.cc/.h。这里承担 pass 管道:layout inference、layout propagation、tensorization、TMA / MMA / warp-specialization 注入、Lower TIRX → TIR/LLVM IR、autotune 搜索、pass 计时与 IR dump。tilelang/autotuner 与 tilelang/contrib(含 contrib/triton、contrib/dsl)则提供高层调度策略与第三方 DSL 桥接。 再往下是 执行层 / 后端层,由 tilelang/backend 与 src/backend 暴露统一的 Backend 注册表,把编译产物分发给各硬件后端:CUDA(tilelang/cuda、src/cuda,含 NVRTC、CuTe DSL、SM75/90/100/120 tensor core、TMA、warp specialization)、ROCm/HIP(tilelang/rocm、src/rocm)、Metal(tilelang/metal、src/metal)、Ascend NPU(tilelang/ascend、src/ascend)、LLVM CPU(src/cpu)、WebGPU(src/webgpu)以及 src/tl_templates 里的模板 kernel。 最底层是 硬件抽象与运行时层,由 src/runtime、3rdparty/tvm、3rdparty/cutlass、3rdparty/hip-headers 提供 TVM runtime FFI、CuTe CUTLASS 模板、HIP 头文件与 CUDA stub。工具层(tilelang/profiler、tilelang/instrumentation、tilelang/tools、benchmark/、maint/)贯穿各层,负责 profiling、layout 可视化、autodd、IR trace / pass diff、正确性评估等开发者体验。

接入层调度层 / 编译器后端层运行时层工具层其他tilelang.__init__ → tilelang.jit:请求请求tilelang.language → tilelang.ir:TIRX IRTIRX IRtilelang.jit → tilelang.engine:LowerLowertilelang.engine → tilelang.transform:pass 管道tilelang.transform → tilelang.backend:后端后端tilelang.backend → tilelang.cuda:CUDACUDAtilelang.backend → tilelang.rocm:ROCmROCmtilelang.backend → tilelang.metal:MetalMetaltilelang.backend → tilelang.ascend:AscendAscendtilelang.cuda → 3rdparty.cutlass:CuTe DSLCuTe DSLtilelang.rocm → 3rdparty.hip-headers:HIP 头HIP 头tilelang.jit → tilelang.cache:缓存缓存tilelang.profiler → tilelang.jit:profileprofiletilelang.autotuner → tilelang.engine:tunetunetilelang.engine → src.runtime:launchlaunchtilelang.transform → tilelang.layout:layoutlayouttilelang.tools → tilelang.transform:调试 IR调试 IRtilelang.__init__tilelang.__init__tilelang.languagetilelang.languagetilelang.jittilelang.jittilelang.enginetilelang.enginetilelang.transformtilelang.transformtilelang.autotunertilelang.autotunertilelang.backendtilelang.backendtilelang.contribtilelang.contribtilelang.cudatilelang.cudatilelang.rocmtilelang.rocmtilelang.metaltilelang.metaltilelang.ascendtilelang.ascendsrc.runtimesrc.runtimetilelang.toolstilelang.toolstilelang.layouttilelang.layouttilelang.profilertilelang.profilertilelang.instrumentationtilelang.instrumentationtilelang.cachetilelang.cachetilelang.analysistilelang.analysistilelang.carvertilelang.carvertilelang.irtilelang.irsrc.transformsrc.transformsrc.backendsrc.backend3rdparty.tvm3rdparty.tvm3rdparty.cutlass3rdparty.cutlass3rdparty.hip-headers3rdparty.hip-headers
TileLang 五层架构:接入层(DSL + JIT)→ 调度层(Engine + Transform + Backend)→ 后端层(CUDA/ROCm/Metal/Ascend/CPU/WebGPU)→ 运行时层(Runtime launch)→ 工具层(Profiler/Layout/Tools)
模块 / 路径职责 · 入口 · 依赖
tilelang.__init__
tilelang/__init__.py
1 文件
顶层入口,聚合 compile / jit / dsl 导出
入口:compile, jit, dsl, language, TensorSupplyType, PassConfigKey
依赖:tilelang.jit, tilelang.ir, tilelang.engine, tilelang.layout
用户第一次 import tilelang 的入口;README 与 pyproject 都把它作为主包门面
tilelang.jit
tilelang/jit
约 6 文件(目录)
JIT 编译与缓存,把 DSL 函数编译成可调用 kernel
入口:JITKernel, JITFunction, LowerAndImportEngine
依赖:tilelang.ir, tilelang.engine, tilelang.cache, tilelang.backend
提交热点之一(6 次);负责 Lower TIRX → 后端、缓存命中、输出索引保留(#3318)
tilelang.language
tilelang/language
约 8 文件
Pythonic DSL 原语:Kernel / Tensor / Pipelined / copy / gemm 等
入口:T.Kernel, T.Tensor, T.Pipelined, T.copy, T.gemm, T.alloc_shared, T.alloc_fragment
依赖:tilelang.ir, tilelang.analysis, tilelang.layout
README 与 docs 自称的 Pythonic 语法层;最近一次提交涉及 float source bitwise reduce 修复(#3332)
tilelang.ir
tilelang/ir.py
约 1 文件(+ C++ 侧 ir.cc/.h)
TIRX IR 定义与操作入口,衔接 TVM Script
入口:TIRX IR node wrappers, tir.Buffer, tir.PrimFunc
依赖:src/ir.cc, src/ir.h, 3rdparty/tvm
README 明确 TIRX 是基于 TVM TIR 发展出的新 IR;#2216 完成 TIRX migration
tilelang.transform
tilelang/transform
约 15 文件
Pass 管道:layout inference、tensorization、TMA/MMA 注入、warp specialization
入口:PassContext, pass injection APIs, LayoutInferPass, TensorizePass
依赖:tilelang.ir, tilelang.analysis, tilelang.layout, src/transform
编译器核心;#2622 pass 计时、#2449 pass visualizer、#2725 IR Lower Trace 都围绕它
tilelang.engine
tilelang/engine
约 3 文件
高层编译引擎:调度 pass、codegen、profiler 绑定
入口:Engine, LowerModule, Codegen
依赖:tilelang.transform, tilelang.backend, tilelang.ir
把 transform 与 backend 串起来;#2441 tile scheduler、#2442 backend registry 都落在这里
tilelang.backend
tilelang/backend
约 3 文件
后端注册表,把编译产物分发给 CUDA/ROCm/Metal/Ascend/CPU/WebGPU
入口:Backend, BackendRegistry
依赖:src/backend, tilelang.cuda, tilelang.rocm, tilelang.metal, tilelang.ascend, tilelang.cpu, tilelang.webgpu
#2442 把 device/host CodeGen 移到 registry 后面;README 列出的多后端都在此调度
tilelang.cuda
tilelang/cuda
约 20 文件
CUDA 后端:NVRTC、CuTe DSL、SM75/90/100/120 tensor core、TMA、warp specialization
入口:CUDASourceCodeKernel, CUDACodegen, TensorCoreConfig
依赖:src/cuda, 3rdparty/cutlass
提交热点(3 次);README 详列 SM120 / MXFP8 / NVF4 / TMA / cluster copy / warp spec
tilelang.rocm
tilelang/rocm
约 8 文件
ROCm/HIP 后端:CDNA4、RDNA3/RDNA4、WMMA
入口:ROCmCodegen, HIPBackend
依赖:src/rocm, 3rdparty/hip-headers
#2044 RDNA3 WMMA、#2132 CDNA4 MXFP4;3rdparty/hip-headers 提供 HIP 头
tilelang.metal
tilelang/metal
约 5 文件
Metal 后端:simdgroup、cooperative tensor GEMM
入口:MetalCodegen, MetalBackend
依赖:src/metal
#1869 simdgroup GEMM、#2252 Metal 4 cooperative tensor;Apple Silicon 可直接用
tilelang.ascend
tilelang/ascend
约 6 文件
Ascend 950 NPU 官方后端:PTODSL、SIMT/SIMD 向量编程
入口:AscendBackend, PTODSLCodegen
依赖:src/ascend
README 2026-09-30 公告;#3469 提交 vendored PTODSL SIMT allreduce helpers
tilelang.autotuner
tilelang/autotuner
约 6 文件
自动调优:并行 autotune、eager-mode、pass configs 搜索
入口:Autotuner, PassConfigKey search
依赖:tilelang.engine, tilelang.jit
提交热点(2 次);#2159 并行 autotune、#1906 eager-mode、#3393 结构化 trial 报告
tilelang.contrib
tilelang/contrib
约 5 文件
第三方 DSL 桥接与扩展(triton、dsl 等)
入口:contrib.triton, contrib.dsl
依赖:tilelang.language, tilelang.backend
提交热点(5 次);README 提到 multi-backend dialect 与第三方桥接
src.transform
src/transform
约 30 文件
C++ 侧 pass 实现:layout、tensorization、TMA/MMA 注入
入口:C++ pass classes, LowerTIRX
依赖:src/ir.cc, 3rdparty/tvm
与 tilelang/transform 镜像;#3427 For loop step 修复、#3245 GEMM dtype 拒绝
src.backend
src/backend
约 10 文件
C++ 侧后端注册表与 codegen 分发
入口:BackendRegistry, CodegenDevice, CodegenHost
依赖:src/cuda, src/rocm, src/metal, src/ascend, src/cpu, src/webgpu
#2442 把 codegen 移到 registry 后;与 tilelang/backend 对应
tilelang.profiler
tilelang/profiler
约 3 文件
Profiler:kernel 计时、tensor supply、正确性验证
入口:Profiler, assert_allclose
依赖:tilelang.jit, tilelang.backend
maint/gemm/correctness_evaluation.py 用它做正确性评估
tilelang.cache
tilelang/cache
约 2 文件
编译缓存:跨主机 CUDA binary cache、TIRX 缓存
入口:CacheManager
依赖:tilelang.jit
#2459 跨主机 CUDA binary cache 落在这里
tilelang.layout
tilelang/layout
约 5 文件
Layout 表示、推理与可视化
入口:Layout, LayoutInference
依赖:tilelang.analysis, tilelang.ir
#1811 布局可视化改进;maint/layout_inference 与之配套
tilelang.analysis
tilelang/analysis
约 6 文件
静态分析:循环嵌套、指针、数据类型检查
入口:LoopNestAnalysis, DTypeAnalysis
依赖:tilelang.ir
.agents/skills/tilelang-semantic/scripts/scan_loop_nesting.py 体现其能力
tilelang.carver
tilelang/carver
约 2 文件
Kernel carving / 切片,用于 autotune 与调试
入口:Carver
依赖:tilelang.ir, tilelang.transform
README 未详述;与 autotune 和 pass 配合
tilelang.instrumentation
tilelang/instrumentation
约 2 文件
IKET profiler 集成、CUDA timeline instrumentation
入口:IKETProfiler
依赖:tilelang.profiler
#2515 IKET profiler 集成;README 2026-07-12 提及
tilelang.tools
tilelang/tools
约 8 文件
开发者工具:compile_only、analyzer、layout_visualization、autodd、lower_trace、pass_diff、iket
入口:compile_only, analyzer, layout_visualization, autodd, lower_trace, pass_diff, iket
依赖:tilelang.jit, tilelang.transform, tilelang.ir
docs/tools 列出这些工具;#2725 lower_trace、#2375 pass_diff、#2045 autodd frozen regions
src.runtime
src/runtime
约 15 文件
C++ 运行时:kernel launch、stream、event、CUDA stub
入口:Runtime, LaunchKernel
依赖:3rdparty/tvm
#1821 CUDA runtime stubs 在此;支撑 JIT 执行
3rdparty.tvm
3rdparty/tvm
大型子模块
TVM 核心:TIR、TIRX、FFI、runtime
入口:tvm.script, tvm.tir, tvm.runtime
依赖:无(底层)
README 明确 TVM 是基础;TIRX 自此发展
3rdparty.cutlass
3rdparty/cutlass
大型子模块
CUTLASS / CuTe DSL 模板库
入口:CudaF, GemMma
依赖:无(底层)
CUDA CuTe DSL 后端依赖;#1421、#1855、#3414 都围绕它
3rdparty.hip-headers
3rdparty/hip-headers
约 40 文件
HIP 头文件,支持 ROCm 后端
入口:hip/hip_runtime.h, hip/amd_detail/*
依赖:无(底层)
目录树可见完整 HIP 头;与 tilelang/rocm 配套
目录树(按文件数)
  • testing/ 499 个文件
    .gitkeep, ascend, conftest.py, cpp, python
  • tilelang/ 414 个文件
    __init__.py, _ffi_api.py, _typing.py, analysis, ascend, autodd.py, autotuner, backend, cache, carver, contrib, cpu, cuda, dtypes.py
  • src/ 398 个文件
    ascend, backend, config.h, cpu, cuda, ir.cc, ir.h, layout, metal, op, rocm, runtime, span_utils.cc, span_utils.h
  • examples/ 387 个文件
    amd, analyze, ascend, attention_sink, autodd, aws, bitnet-1.58b, block_causal_attention, blockscaled_gemm_sm100, blocksparse_attention, blocksparse_gemm, cast, conftest.py, convolution
  • maint/ 75 个文件
    gemm, host_checks, layout_inference, precision, scripts
  • docs/ 62 个文件
    .gitignore, CNAME, Makefile, README.md, _static, compiler_internals, conf.py, deeplearning_operators, developer_guide, get_started, index.md, make.bat, privacy.md, programming_guides
  • 3rdparty/ 40 个文件
    .gitignore, cutlass, hip-headers, tvm
  • .agents/ 21 个文件
    skills
  • benchmark/ 19 个文件
    blocksparse_attention, compile_speed, mamba2, matmul, matmul_fp8, matmul_metal
  • .github/ 11 个文件
    ISSUE_TEMPLATE, dependabot.yml, workflows
  • docker/ 10 个文件
    Dockerfile.cu118, Dockerfile.cu120, Dockerfile.cu121, Dockerfile.cu123, Dockerfile.cu124, Dockerfile.cu125, Dockerfile.cu126, Dockerfile.cu128, Dockerfile.rocm, README.md
  • images/ 10 个文件
    MatmulExample.png, MatmulExample.svg, ascend-950-backend-architecture.png, ascend-950-throughput.png, logo-row.svg, mha_performance_h100.png, op_benchmark_a100_wq_gemv.png, op_benchmark_consistent_gemm_fp16.png, op_benchmark_h100.png, op_benchmark_mi300_fp16_gemm_normalized_latency.png
  • .claude/ 9 个文件
    skills
  • cmake/ 7 个文件
    BackendOptions.cmake, FindPipCUDAToolkit.cmake, ccache_strip_pep517.py, find_pip_cuda.py, generate_windows_import_lib.py, load_tvm.cmake, pypi-z3
  • .clang-format/ 1 个文件
  • .coderabbit.yaml/ 1 个文件

一次调用怎么流过这些模块

一次 TileLang kernel 请求从用户代码开始。用户在 Python 中用 T.Kernel、T.Tensor、T.Pipelined、T.copy、T.gemm 等 DSL 原语(tilelang/language)描述 kernel;这些原语被 TVM Script 解析器捕获成 TIRX IR(tilelang/ir.py + src/ir.cc/.h),缓存在 tilelang/cache。 当用户调用 tilelang.compile(fn) 或用 @tilelang.jit 装饰时,tilelang/jit 的 JITFunction 接管:先查缓存,未命中则进入 tilelang/engine 的 Engine 调度。Engine 把 TIRX 喂进 tilelang/transform 与 src/transform 的 pass 管道——依次做 layout inference、layout propagation、tensorization、TMA/MMA/warp-specialization 注入、Lower TIRX → TIR/LLVM IR。tilelang/autotuner 可在这一层搜索 pass configs 与 tile 尺寸(#2159 并行 autotune、#1906 eager-mode)。 Lower 后的 IR 进入 tilelang/backend + src/backend 的 BackendRegistry,按目标硬件分发给具体后端:CUDA(tilelang/cuda、src/cuda)走 NVRTC 或 CuTe DSL(3rdparty/cutlass),ROCm(tilelang/rocm、src/rocm)走 HIP,Metal(tilelang/metal、src/metal)走 Metal Shading Language,Ascend(tilelang/ascend、src/ascend)走 PTODSL,CPU(src/cpu)走 LLVM IR,WebGPU(src/webgpu)走 WGSL。后端 codegen 产出可执行二进制或源码。 tilelang/jit 把产物包装成 JITKernel 返回给用户;用户像调 PyTorch op 一样传入 tensor,由 src/runtime 的 runtime 在指定 stream 上 launch kernel。tilelang/profiler 的 Profiler 可包裹 JITKernel,提供计时与 assert_allclose 正确性验证(maint/gemm/correctness_evaluation.py 大量使用)。tilelang/instrumentation 的 IKET profiler 可插入 CUDA timeline 做细粒度分析。tilelang/tools 的 lower_trace、pass_diff、layout_visualization、autodd 则用于调试 IR 变化与定位性能瓶颈。 关键调度点:① tilelang/jit 的缓存与 Lower 触发;② tilelang/engine 的 pass 管道调度;③ tilelang/backend 的 BackendRegistry 分发;④ tilelang/autotuner 的 trial 调度;⑤ src/runtime 的 kernel launch。关键数据结构:TIRX IR(src/ir.cc/.h)、JITKernel(tilelang/jit)、Backend 注册表(tilelang/backend)、PassContext(tilelang/transform)、Profiler(tilelang/profiler)。

1用户写 DSLT.Kernel / T.Pipelined / T.gemm · tilelang/language2解析成 TIRXTVM Script parser · tilelang/ir.py, src/ir.cc3JIT 编译触发JITFunction / JITKernel · tilelang/jit4Engine 调度 passEngine + PassContext · tilelang/engine, tilelang/transform5Backend 分发BackendRegistry · tilelang/backend, src/backend6CUDA codegenCUDACodegen / CuTe DSL · tilelang/cuda, 3rdparty/cutlass7Runtime launchRuntime · src/runtime8Profiler 验证Profiler.assert_allclose · tilelang/profiler9结果返回JITKernel.__call__ · tilelang/jit

关键类型与函数

名称路径用途
JITKernel / JITFunctiontilelang/jitJIT 编译后的可调用 kernel 对象,用户最终调用它
T.Kernel / T.Tensor / T.Pipelinedtilelang/languagePythonic DSL 原语,用户写 kernel 的入口
TIRX IRtilelang/ir.py, src/ir.cc, src/ir.hTileLang 基于 TVM TIR 发展的中间表示,贯穿编译管道
Enginetilelang/engine高层编译引擎,调度 pass 与 codegen
Backend / BackendRegistrytilelang/backend, src/backend后端注册表,按硬件分发编译产物
PassContext / PassConfigKeytilelang/transformPass 管道上下文与配置键,控制优化行为
Layout / LayoutInferencetilelang/layoutLayout 表示与推理,驱动 tensorization 与 TMA
Profilertilelang/profilerKernel 计时与正确性验证(assert_allclose)
Autotunertilelang/autotuner自动调优 pass configs 与 tile 尺寸
CUDASourceCodeKerneltilelang/cuda嵌入自定义 CUDA source 的 kernel(#1970)
PTODSLCodegen / AscendBackendtilelang/ascend, src/ascendAscend 950 NPU 官方后端代码生成
Carvertilelang/carverKernel carving / 切片,用于 autotune 与调试
IKETProfilertilelang/instrumentationCUDA timeline instrumentation 与 profiling(#2515)

扩展点

最近在动的地方

建议阅读顺序

  1. README.md(项目定位、最新 News、后端列表)
  2. tilelang/__init__.py(顶层入口与导出)
  3. tilelang/language(DSL 原语:T.Kernel / T.Tensor / T.Pipelined / T.gemm)
  4. examples/gemm/example_gemm.py(最简 GEMM 示例,看 DSL 如何写)
  5. tilelang/jit(JIT 编译、缓存、Lower 触发)
  6. tilelang/ir.py + src/ir.cc/.h(TIRX IR 定义)
  7. tilelang/transform(pass 管道:layout、tensorization、TMA/MMA 注入)
  8. tilelang/engine(Engine 调度 pass 与 codegen)
  9. tilelang/backend(BackendRegistry 分发)
  10. tilelang/cuda(CUDA 后端:NVRTC、CuTe DSL、tensor core、TMA)
  11. tilelang/autotuner(自动调优与 pass configs)
  12. src/runtime(kernel launch 与 CUDA stub)
  13. 3rdparty/tvm、3rdparty/cutlass、3rdparty/hip-headers(底层依赖)

三、本地跑起来(没有 GPU 的 Mac)

安装

  1. # 1. 克隆仓库(带子模块:3rdparty/tvm、cutlass、hip-headers)
  2. git clone --recurse-submodules git@github.com:tile-ai/tilelang.git
  3. cd tilelang
  4. # 2. 创建虚拟环境(Apple Silicon 推荐 uv,也可 python3 -m venv .venv)
  5. uv venv --seed .venv
  6. source .venv/bin/activate
  7. python3 -m pip install --upgrade pip setuptools wheel "build[uv]"
  8. # 3. 安装开发依赖(含测试、lint 等 extras)
  9. uv pip install --requirements requirements-dev.txt
  10. # 4. 可编辑安装 tilelang(无需 host CUDA toolkit,无 GPU 也能构建)
  11. python3 -m pip install --no-build-isolation --verbose --editable .
  12. # 说明:pyproject.toml 的 [project.optional-dependencies] nvcc 组(nvidia-cuda-nvcc 等)必须跳过,否则 CMake 会尝试 FindCUDAToolkit 并失败。
  13. # 若后续想尝试 Metal,确保 macOS 13+ 且 Xcode CLI tools 安装;Metal backend 由 src/metal + tilelang/metal 提供,无需额外 pip 包。
  14. # 若只想读代码、跑 CPU 测试,可跳过步骤 4,改用 PYTHONPATH=src python3 ... 直接运行(见 CONTRIBUTING 的 Working from Source 建议)。

哪些路径能真跑

['可真正执行(Apple Silicon / CPU):', '- CPU 后端:src/cpu/、tilelang/cpu/,通过 LLVM 后端 lowering 并执行(README 2026-03-10 条目 "CPU T.gemm")。', '- Metal 后端:src/metal/、tilelang/metal/,含 simdgroup GEMM 与 cooperative-tensor GEMM(README 2026-05-22、2026-07-28);需实际 M 系列 GPU 才能跑,MacBook Air 可编译但只能验证编译产物。', '- LLVM CPU 测试子集:testing/python 中部分通用测试 + testing/cpp 中与 backend 无关的 IR/transform 测试。', '- examples/ 中少量 CPU 友好示例:examples/elementwise/、examples/gemm/(部分带 target=cpu 或自动后端选择)。', '只能读代码或需 Colab T4 验证:', '- tilelang/cuda/、src/cuda/、3rdparty/cutlass、3rdparty/hip-headers —— 任何含 T.Kernel + T.gemm/T.copy/T.mma_* 的 CUDA 示例均需 GPU。', '- examples/flash_attention/、examples/deepseek_mla/、examples/dequantize_gemm/、examples/gemm_fp8/、examples/blocksparse_attention/ 等依赖 tensor core / TMA / warp specialization 的示例。', '- benchmark/matmul、benchmark/blocksparse_attention 等需要实际 GPU 测时的基准。', '- Ascend NPU(tilelang/ascend、src/ascend、testing/ascend)需华为硬件。', '- ROCm(tilelang/rocm、src/rocm)需 AMD GPU。']

最小可运行

  1. # 1. 最小 elementwise 编译 + 执行(CPU 后端,验证 tilelang.compile 与 JIT 整条链路)
  2. PYTHONPATH=src python3 -c "import tilelang; from tilelang import language as T; print(tilelang.__version__)"
  3. # 2. CPU 端到端 GEMM 正确性(LLVM 后端),参考 README 2026-03-10 条目
  4. python3 examples/gemm/example_gemm.py # 若脚本默认 target=cuda,需手动改 target=cpu 或加 --target cpu(需验证脚本参数)
  5. # 3. 在 Colab T4 上跑最小 FlashAttention 示例(最终确认 CUDA 后端可用)
  6. # 上传仓库后执行:
  7. python3 examples/flash_attention/example_mha_fwd_bhsd.py
  8. # 4. 编译-only 验证(不执行,仅检查 IR lowering 是否通过)
  9. python3 -c "from tilelang import compile; compile(lambda: None, out_idx=[])" # 占位,实际需替换为最小 prim_func
  10. # 5. Metal 后端编译验证(M 系列 Mac,不执行)
  11. python3 -c "from tilelang.metal import is_metal_available; print(is_metal_available())"
  12. # 6. 运行 CPU 子集测试(见 test_suite)

测试

['框架:pytest(根目录 testing/),部分示例自带 test_*.py 通过 pytest 发现。', '目录结构:testing/python/(主要 Python 测试)、testing/cpp/(C++ 单测)、testing/ascend/(Ascend 专用)。', '只跑 CPU 子集(MacBook Air 无 GPU):', ' python3 -m pytest testing/python -k "cpu or metal or elementwise or gemm" --timeout=120 -x', ' # 或更保守:只跑与 CUDA 无关的通用测试', ' python3 -m pytest testing/python -k "not cuda and not rocm and not ascend and not hip" --timeout=120', '示例自带测试(examples/ 下 test_*.py):', ' python3 -m pytest examples/elementwise/test_example_elementwise.py examples/gemv/test_example_gemv.py -x', '耗时估计:CPU 子集(约几十个测试)在 Apple Silicon 16GB 上约 5–15 分钟;完整 testing/ 含 GPU 测试会因 skip 而很快,但含编译的测试可能较慢。', '注意:testing/conftest.py 可能自动检测 GPU 并 skip 无设备测试,需检查其 skip 标记逻辑。']

调试

CI

GitHub Actions(.github/workflows/),多后端矩阵:CUDA(cu118/cu120/.../cu128)、ROCm、Metal、Ascend、CPU/LLVM;PR 会跑 lint(pre-commit: ruff + clang-format + pymarkdown)、单元测试、示例回归测试。无 GPU 贡献者提交会被 GPU 测试 skip,但 lint 与 CPU 测试必须通过;CodeRabbit 自动评审(最近提交 fd2ddda 启用)。具体 job 名与阈值需查看 .github/workflows/*.yml 确认。

坑

四、维护者与社区

Release 节奏约为每月一版(v0.1.13 2026-08-03 → v0.1.14 2026-09-02 → v0.1.15 2026-09-30),提交非常活跃,最近提交集中在 2026-10-10 至 10-12 之间,涉及 SM120、CuTeDSL、JIT、Autotune、Ascend、Transform 等多个子系统,说明处于多后端扩展与稳定性并重的密集迭代期。提交热点(30 天):testing/python (20)、tilelang/jit (6)、tilelang/contrib (5)、src/cuda (3)、tilelang/autotuner (2)、testing/ascend (2)。

谁角色依据
Lei Wang (LeiWang1999)项目创建者与主要维护者,邮件 leiwang1999@outlook.com,在 pyproject.toml 的 maintainers 字段列出;在 Issue #2986 被 assignee,负责核心方向决策。pyproject.toml maintainers 字段;Issue #2986 assignee 为 LeiWang1999
Dexterai活跃贡献者,负责 transform / autotuner / CUDA 相关工作,被 assign 到 Issue #3120。Issue #3120 assignees 包含 Dexterai
CodeRabbit (bot)自动化代码评审机器人,已被启用 reviewer auto-assignment。Commit fd2ddda [Maint] Enable CodeRabbit reviewer auto-assignment (#3487);文件 .coderabbit.yaml 存在
社区贡献者群体大量 PR 来自不同作者(Ascend 后端、BugFix、Transform、Autotune、Profiler、Metal、ROCm 等方向),说明有多位核心与外围维护者共同推进。最近 PR 涉及 #3488、#3474、#3476、#3328、#3486、#3481、#3417、#3350、#3330、#3452 等多个不同方向与作者

流程与 Review 风格

贡献流程:1) Fork 仓库并用 SSH 克隆(需 --recurse-submodules,含 3rdparty/tvm、cutlass、hip-headers);2) 用 uv venv --seed .venv 创建环境并安装 requirements-dev.txt;3) 运行 pre-commit install --install-hooks 安装钩子;4) 用 python3 -m pip install --no-build-isolation --verbose --editable . 可编辑安装;5) 提交 PR 前运行 bash format.sh --files <changed-file> 格式化,Python 由 Ruff 检查,C++ 由 clang-format 检查,C++ API 风格需符合 docs/developer_guide/cpp_style.md;可用 python3 maint/scripts/audit_cpp_api_style.py 做命名与头文件边界审计;6) 测试用 python3 -m pytest testing;7) 搜索已有 Issue 再新建,新建 Issue 必须包含最小复现代码与清晰描述;8) 可在 Issues 中提问;9) 带 [♥ help] 标签的 Issue 是入门入口。无 CLA/DCO 证据,pre-commit 是主要门禁。

Review 较密集且具体:BugFix 类 PR(如 #3414、#3427、#3318、#3300、#3400、#3452)集中在 JIT/CUDA/Transform 子系统的正确性边界,说明维护者对 warp specialization、TMA、NVRTC、shared memory race 等底层细节审查严格;Feature 类 PR(如 #3476 autotune 并行、#3393 structured trial reporting、#3481 实验性 MIR)接受方向性扩展;同时启用 CodeRabbit 自动化评审(#3487),表明对代码风格与常见模式有一层自动把关。Issue 响应较快,如 #3301 有 12 条讨论,多个 bug 在数日内就有对应 PR。

渠道

这里的规矩

维护者现在最想要的帮助

五、切入方案

建议长期负责:CPU / Metal / WebGPU 后端的 correctness 测试基础设施与调试工具链(transform pass 可视化、IR dump、JIT 缓存可观测性)
这三条路径都可在 MacBook Air 上完整跑通(LLVM CPU 后端可执行、Metal 后端可编译、WebGPU 可编译-only),且与候选人 GPU kernel 经验高度相关——理解 lowering、tensorization、layout inference 后,自然能产出高质量的 IR 级测试与调试工具;同时这些模块是 CUDA/Ascend 后端的“公共基础设施”,做深之后可顺理成章地扩展到其他后端,通向核心维护者路径。

它现在缺什么(你无 GPU 也能补)

缺口依据为什么是你
CPU / Metal 后端缺乏系统性的 correctness 基准与回归测试套件maint/gemm 目录下的 correctness_evaluation*.py 全部面向 CUDA(含 SM70、tcgen05、SM120、NVF4 等 GPU 特化路径),而 CPU 后端(src/cpu、tilelang/cpu)和 Metal 后端(src/metal、tilelang/metal)在 testing/ 与 maint/ 中几乎没有对应的 correctness 评估脚本;README 2026-03-10 仅提到“CPU T.gemm”已支持,README 2026-05-22、2026-07-28 提到 Metal simdgroup / cooperative-tensor GEMM,但未见公开的 CPU/Metal 精度验证工具。候选人擅长算子精度验证与 PyTorch 对照,且 MacBook Air 可直接跑 CPU 后端(LLVM)与 Metal 后端编译产物,无需任何 GPU 即可产出可重复的 correctness 测试。
Metal 后端缺少从 DSL 高层语义到 simdgroup / cooperative-tensor 的端到端示例与文档docs/programming_guides 与 docs/deeplearning_operators 主要覆盖 CUDA 路径(matmul、matmul_sparse、deepseek_mla 等),Metal 只在 compiler_internals/metal_tilelang_development 有一篇内部开发指南;benchmark/matmul_metal 仅一个 benchmark_matmul_metal.py,无对应的回归测试或 step-by-step 教程。候选人在 Apple Silicon 上可直接编译、调试 Metal pipeline,能写出“无 GPU 也能跑”的 Metal 教程与示例,填补当前文档空白。
Ascend 950 后端的算子覆盖与测试基础设施仍不完整提交热点中 testing/ascend 出现 2 次、tilelang/ascend 与 src/ascend 各 1 次;examples/ascend 含 GEMM、FlashAttention、RMSNorm、vecadd 等,但缺少 SparseMLA、MoE、Dequant GEMM 等 CUDA 侧已丰富的算子;maint/ 下完全没有 ascend 子目录。候选人熟悉 CUDA/Triton 算子实现,可在无 GPU 条件下把 CUDA 侧已有算子“移植”为 Ascend 版本或至少补全测试基础设施(对照 PyTorch / torch_npu),这是当前提交热点方向。
transform pass 的 IR dump / pass diff / pass timing 工具链缺少文档与示例README 2026-07-22、2026-07-16、2026-06-16 分别引入 IR Lower Trace、compiler pass timing、Pass Diff、Pass Visualizer;docs/tools 仅列 compile_only、analyzer、layout_visualization、autodd、lower_trace、pass_diff 六个工具,无使用示例或端到端教程。候选人在编译器诊断与 kernel 调试上有经验,可在无 GPU 条件下跑通这些 pass 调试工具(它们作用于 IR,不依赖具体 GPU),并产出文档与示例。
JIT 缓存、kernel 缓存与跨主机 CUDA 二进制缓存缺少可观测性README 2026-06-26 引入跨主机 CUDA 二进制缓存;tilelang/jit 在 30 天提交热点中排第 2(6 次提交);tilelang/cache 目录存在但 README 与 docs 中几乎未解释缓存目录结构、命中统计、失效策略。候选人熟悉训练系统性能分析,可在 CPU/Mac 上设计缓存命中率统计、调试接口与文档,不依赖 GPU。
WebGPU 后端仍处于实验状态,缺测试与示例src/webgpu 存在于目录树,但 examples/、testing/、maint/ 下未见 webgpu 子目录;README 未在 news 中正式宣传 WebGPU,仅目录树暗示其存在。候选人可在 Apple Silicon 上通过 Metal 或 LLVM 路径间接验证 WebGPU 的编译流程,或至少补全编译-only 的测试用例,无需真实 GPU。
autotuner 的并行调优与结构化试验报告刚落地,缺少使用文档与真实 benchmark 对比README 2026-05-11 引入 parallel autotuning,2026-10-12 提交“[Feature][Autotune] Add opt-in structured trial reporting (#3393)”;docs/programming_guides/autotuning 存在,但 benchmark/ 下只有 matmul、matmul_fp8、matmul_metal、mamba2、blocksparse_attention、compile_speed 六个子目录,无 autotune 结果归档或复现脚本。候选人可在 CPU / Metal 后端上跑 autotune 流程(autotune 本身支持 LLVM 后端),产出可复现的调优报告与文档,验证 #3393 的结构化 trial 输出。
host_checks 静态检查器覆盖范围有限maint/host_checks/ 仅有 01_num_args_mismatch.py、02_pointer_type_error.py、03_ndim_mismatch.py 三个脚本,远少于实际可能出现的用户错误模式(layout 推断失败、dtype 不匹配、scope 错误等)。候选人在 CUDA/Triton/PyTorch 内核开发中积累了大量“常见错误模式”,可在无 GPU 条件下为 host_checks 新增静态检查用例,直接降低入门门槛。

第 1–30 天:看懂并露面

第 31–60 天:稳定产出

第 61–90 天:接管一块

第一批 PR

题目范围为什么安全
[Test][CPU] Add correctness evaluation for GEMM / GEMV / elementwise on LLVM CPU backend在 maint/ 下新增 maint/cpu/correctness_evaluation_*.py,参照 maint/gemm/correctness_evaluation.py 的 PyTorch 对照模式,覆盖 CPU 后端已支持的算子。只新增测试文件,不改 src/ 或 tilelang/ 核心代码;全部在 MacBook Air 的 LLVM 后端上可验证,无需 GPU。
[Test][Metal] Add Metal compile-only correctness evaluation for GEMM在 maint/ 下新增 maint/metal/correctness_evaluation_metal.py,验证 Metal 后端能编译出正确形状的 kernel 产物(可对比 CPU 后端输出做 bit-exact 或近似对照)。只新增测试与文档,不修改 Metal codegen 本身;编译失败或精度偏差都会以测试失败形式暴露,不会破坏现有功能。
[Docs][Tools] Add end-to-end examples for lower_trace and pass_diff在 docs/tools/ 下新增 lower_trace.md 与 pass_diff.md,配合 examples/ 中已有的一个简单 kernel(如 examples/gemm/example_gemm.py)展示 IR 在每一步 pass 的变化。纯文档与示例,不修改源码;可在 CPU 后端上完整复现,无需 GPU。
[Maint][HostChecks] Add static checks for dtype mismatch and scope errors在 maint/host_checks/ 新增 04_dtype_mismatch.py、05_scope_error.py,参考已有 01/02/03 的 pattern。只新增静态检查脚本,不修改核心编译器;失败模式是“检查到错误时抛异常”,不会误杀合法代码(因为只针对明确的非法模式)。
[Test][Autotune] Add CPU backend autotune reproducer with structured trial reporting在 benchmark/ 下新增 benchmark/autotune_cpu/,使用 #3393 引入的结构化 trial 报告功能,在 CPU 后端上跑通 autotune 并输出可复现的 trial 日志。只新增 benchmark 脚本,不修改 autotuner 核心;失败只会是 autotune 未收敛或日志格式异常,不影响现有功能。

怎么知道自己站住了

风险与对策
  • CPU 后端支持的算子覆盖面比 CUDA 窄,写 CPU correctness 测试时可能遇到“算子已支持但精度/shape 边界未定义”的情况。:先以 README 与 examples/ 中已明确宣传的算子(GEMM、GEMV、elementwise)为范围,遇到边界 case 先提 issue 而非直接改核心代码。
  • Metal 后端在 macOS 上编译可能受 Xcode 版本、Metal 版本影响,导致编译失败误判为后端 bug。:在 CI 或本地先记录 metal 编译器版本与产物哈希,区分“编译环境变化”与“后端 codegen 回归”。
  • transform pass 回归测试高度依赖 IR 的字符串表示,pass 输出微调会导致测试大面积 flaky。:用结构化比较(如 IR 的 AST 或 pass_diff 工具)而非纯文本 diff;测试只断言关键变换(如 tensorization 是否发生),不要求全文匹配。
  • 项目处于快速迭代期(每月一版),maintenance 脚本与 host_checks 可能与新 API 不同步。:每个 PR 都跑 pre-commit run --all-files 与 python3 -m pytest testing,并在 CI 矩阵中锁定当前主分支版本。
  • 无 GPU 导致无法验证 CUDA/Ascend 后端的真实性能与精度,贡献范围受限。:聚焦 CPU/Metal/测试/文档/tools 这些“后端无关”的基础设施,用 Colab 免费 T4 仅在最终确认时跑一次 CUDA 编译(README 允许 host-toolchain-free builds)。
  • 社区沟通以 Discord 与中文为主,时区与语言可能成为融入障碍。:前 30 天至少发 3 次高质量 Discord 消息(自我介绍、PR review 回复、issue 分析),并主动在 issue 中用英文撰写技术总结。
  • CodeRabbit 自动 review 已启用,PR 若风格不符会被自动打回。:每次提交前本地跑 bash format.sh --files <changed> 与 pre-commit run --all-files,避免低级风格问题。
  • ownership_target 模块(CPU/Metal 测试)在现有维护者眼中优先级可能低于 CUDA/Ascend 新功能,PR 响应慢。:把测试基础设施与“新功能可用”绑定(例如“让 Ascend 新算子合并前必须先过 CPU 对照测试”),提升模块的战略可见度。

六、怎么介入这个项目

社区入口:img.shields.io/badge/Discord-%235865F2.s · discord.gg/TUrHyJnKPG

建议顺序

成为长期维护者的路径
  • 从 maint/host_checks/04_dtype_mismatch.py、maint/cpu/correctness_evaluation_*.py、docs/tools/lower_trace.md 等纯新增文件起步,不触碰 src/ 与 tilelang/ 核心代码
  • 用 tilelang/tools 中已有的 lower_trace / pass_diff / compile_only 等调试工具产出文档与示例,维护者已在 #2725、#2375、#2622 投入明显精力
  • 把 CPU 后端的 correctness 评估与 autotune reproducer 落到 maint/ 与 benchmark/ 下,与现有 maint/gemm 结构对齐
  • 在熟悉流程后,可进一步承担 transform pass 稳定性修复(如 #3427、#3245 类)与 Metal codegen bug(#3292)等需要 Apple Silicon 的任务

七、任务卡

任务 1优先 · low · CPU · 2–3 个晚上

[Maint][HostChecks] 新增 dtype 不匹配与 scope 错误的静态检查脚本

已有 PR #331912 条评论更新 2026-10-11

用到的专长:候选人在 CUDA/Triton/PyTorch 内核开发中积累了大量「常见错误模式」直觉,可直接转化为静态检查用例。

目标:在 maint/host_checks/ 下新增 04_dtype_mismatch.py 与 05_scope_error.py,分别捕捉「dtype 不匹配」与「scope(shared/local/fragment)误用」两类明确非法模式,失败时抛出带位置信息的异常而非裸 IndexError / Check failed。

为什么值得长期做:#3301 是近期维护者最活跃的 batch bug(12 条评论、20+ 个 open PR),host_checks 扩展是其中「降低入门门槛、提前拦住用户错误」这一支的直接延续;现有 01/02/03 脚本已建立清晰模式,新增脚本可直接对齐。

怎么介入:#3301 已有 20+ open PR 在推进,但 host_checks 目录下的扩展尚未被认领;engagement 建议明确说明「只做 04/05 两个新增脚本,不动 01/02/03 与核心编译器」。
第一个 PR 的边界:第一个 PR 只新增 maint/host_checks/04_dtype_mismatch.py、05_scope_error.py 与一份 README;不修改任何 src/ 或 tilelang/ 代码,不改动现有 01/02/03 脚本。
第一步:克隆仓库并阅读 maint/host_checks/01_num_args_mismatch.py、02_pointer_type_error.py、03_ndim_mismatch.py 三份脚本,确认它们如何调用编译器前端、如何断言异常类型与消息。
本机怎么复现 / 验证:在 MacBook 上 pip install -e . 后直接 python3 maint/host_checks/04_dtype_mismatch.py 与 python3 maint/host_checks/05_scope_error.py,无需 GPU;编译路径走 LLVM CPU 后端或纯 host 端检查。
认领留言(英文,可直接贴到 Issue)
I'd like to pick up the host_checks extension under #3301: add maint/host_checks/04_dtype_mismatch.py and 05_scope_error.py, following the pattern of 01/02/03. Plan: construct minimal illegal kernels, assert the exception message names the offending dtype/scope, and add a small README under maint/host_checks/. Should I also add positive (legal) cases that must not raise, or keep these scripts negative-only? Happy to open a separate PR under #3301. Target: a few evenings.
大致实施方案
  • 在 maint/host_checks/04_dtype_mismatch.py 中构造一个 dtype 明显不匹配的最小 kernel(例如把 float16 缓冲区当作 int8 传入 T.copy),触发编译并断言抛出的异常消息里包含 dtype 名称
  • 在 maint/host_checks/05_scope_error.py 中构造一个 scope 误用的最小 kernel(例如在 fragment scope 上调用 T.alloc_shared),触发编译并断言异常消息里包含 scope 名称
  • 每个脚本附一个 if __name__ == '__main__': 入口,便于单独运行
  • 在 maint/host_checks/ 下如有 __init__.py 或 README,补充两行说明;没有则新增一份 README 列出 01–05 的职责
  • 本地运行 python3 maint/host_checks/04_dtype_mismatch.py 与 python3 maint/host_checks/05_scope_error.py,确认它们在预期路径下失败、在合法 kernel 下不误报
可能涉及的目录或文件
  • maint/host_checks/ —— 已有 01/02/03 脚本,直接新增 04/05
  • maint/scripts/ —— 如有共享 helper 可复用
  • docs/developer_guide/ —— 可在文档中补一段 host_checks 说明(可选,非必须)
验收方式
  • python3 maint/host_checks/04_dtype_mismatch.py 在非法 dtype 不匹配 kernel 上抛出含 dtype 名称的异常,在合法 kernel 上不抛
  • python3 maint/host_checks/05_scope_error.py 在非法 scope 误用 kernel 上抛出含 scope 名称的异常,在合法 kernel 上不抛
  • pre-commit run --all-files 通过(新增 Python 文件需满足 Ruff 规则)
  • 确认现有 01/02/03 脚本行为不受影响(如有共用的入口脚本,跑一遍)
开工前问题与风险

向维护者确认

  • host_checks 脚本是否只要求「在非法模式上抛异常」,还是也要求「在合法模式上显式不抛」?即测试的正反两面都要覆盖吗?
  • 异常消息的粒度:是要精确匹配字符串,还是只要包含 dtype/scope 名称即可?
  • 是否需要同步更新某个 CI 配置或文档来把 04/05 纳入必跑列表?

风险

  • 误判合法 kernel 为非法(false positive)——需要在多种合法写法上回归
  • 异常消息措辞与现有 01/02/03 不一致,导致风格偏差——先抄它们的断言模式
任务 2优先 · medium · CPU · 3–4 个晚上

[Test][CPU] 为 LLVM CPU 后端新增 GEMM / GEMV / elementwise correctness 评估

无人认领1 条评论更新 2026-10-11

用到的专长:候选人擅长算子精度验证与 PyTorch 对照,这是该任务的直接映射。

目标:在 maint/ 下新增 maint/cpu/correctness_evaluation_gemm.py(与 GEMV / elementwise 合在一份文件或拆成多份均可),参照 maint/gemm/correctness_evaluation.py 的 PyTorch 对照模式,在 LLVM CPU 后端上跑通并比对结果。

为什么值得长期做:#3000 明确指出 bug 是 host-side 编译期行为(arithmetic simplifier 在 lowering 前崩溃),与 GPU 无关;且 README 2026-03-10 已宣称 CPU T.gemm 可用,但 maint/ 下完全没有 CPU 后端的 correctness 评估脚本。这张卡把「CPU 后端到底对不对」做成可重复的回归测试,直接补上 gaps 中指出的 missing CPU correctness 基础设施。

怎么介入:#3000 本身是「ice-on-valid-code」bug,但本卡不修复 #3000,而是借 #3000 所在的「CPU 后端正确性」主题做基础设施;engagement 建议明确说明「先补测试,后续可用它定位 #3000 类 bug」。
第一个 PR 的边界:第一个 PR 只新增 maint/cpu/correctness_evaluation_gemm.py(与可能的 elementwise),不修改 src/ 或 tilelang/ 核心代码,不触碰 maint/gemm/ 现有文件。
第一步:阅读 maint/gemm/correctness_evaluation*.py 的 PyTorch 对照模式,以及 examples/gemm/ 中支持 target=cpu 或自动后端选择的样例。
本机怎么复现 / 验证:在 MacBook 上 pip install -e . 后 python3 maint/cpu/correctness_evaluation_gemm.py,全部走 LLVM CPU 后端;无需 GPU。
认领留言(英文,可直接贴到 Issue)
I'd like to add a CPU-backend correctness suite under maint/cpu/ (starting with GEMM + one elementwise op), mirroring the PyTorch-reference pattern in maint/gemm/correctness_evaluation.py. All tests run on the LLVM CPU backend on my Mac, no GPU needed. Before I start: which dtypes does the CPU backend currently support (fp32 only, or also fp16/bf16)? And should this be wired into CI immediately or land as a local regression script first? Plan to open a PR in 3–4 evenings.
大致实施方案
  • 在 maint/cpu/correctness_evaluation_gemm.py 中实现「随机初始化 + TileLang CPU kernel + PyTorch reference」的对照测试,覆盖 fp32/fp16/bf16 的 GEMM(若 CPU 后端支持)
  • 至少再覆盖一个 elementwise 算子(如 vecadd / scale),验证 CPU 后端的基础数据流
  • 每个测试用例输出 max abs diff 与 pass/fail,失败时非零退出码
  • 在脚本头部注明需要 pip install -e . 且依赖 PyTorch
  • 本地跑 python3 maint/cpu/correctness_evaluation_gemm.py 全部通过
可能涉及的目录或文件
  • maint/cpu/ —— 新增目录,与 maint/gemm/ 平行
  • maint/gemm/correctness_evaluation.py —— 作为 PyTorch 对照模式的参考实现
  • examples/gemm/ 与 examples/elementwise/ —— 作为 kernel 写法参考
验收方式
  • python3 maint/cpu/correctness_evaluation_gemm.py 在 MacBook 的 LLVM CPU 后端上全部 pass
  • 至少覆盖 GEMM + 1 个 elementwise 算子
  • 与 PyTorch reference 的误差在 dtype 对应的正常范围内(fp32 ~1e-5,fp16 ~1e-3)
  • pre-commit run --all-files 通过
开工前问题与风险

向维护者确认

  • CPU 后端当前支持的 dtype 集合是什么?fp16/bf16 在 LLVM 路径下是否已可用,还是只支持 fp32?
  • maint/ 下是否已有隐藏的 CPU 测试(例如 testing/python 中)需要避免重复?
  • 是否需要把这份脚本接入 CI,还是先作为本地回归工具?

风险

  • CPU 后端对某些 dtype 尚未支持,导致测试被迫跳过——需要维护者确认支持矩阵
  • PyTorch reference 的随机初始化若不一致会导致误报——应固定 seed
任务 3可选 · low · CPU · 2 个晚上

[Docs][Tools] 为 lower_trace 与 pass_diff 新增端到端使用示例

无人认领0 条评论enhancement更新 2026-10-08

用到的专长:候选人在编译器诊断与 kernel 调试上有经验,能准确解读每步 pass 的 IR 变化。

目标:在 docs/tools/ 下新增 lower_trace.md 与 pass_diff.md,配合 examples/ 中已有的简单 kernel(如 examples/gemm/example_gemm.py 或 examples/elementwise 中的一个),展示 IR 在每一步 pass 的变化与 diff 输出。

为什么值得长期做:#3459 是 RFC,但正文反复引用 #3444、legalize_pairwise_extents、CopyLoopsOverSrc 等编译期路径,说明维护者正把 transform pass 调试工具链(lower_trace #2725、pass_diff #2375、pass timing #2622、pass visualizer #2449)推到用户面前。docs/tools 目前只有工具列表而无使用示例,这张卡直接补上「第一步该怎么做」的空白。

怎么介入:#3459 本身是 RFC,本卡不改动 T.copy 语义,而是借该 RFC 暴露出的「用户看不懂 pass 调试工具」问题做文档;engagement 建议明确说明「先出文档示例,不动 T.copy 实现」。
第一个 PR 的边界:第一个 PR 只新增 docs/tools/lower_trace.md、docs/tools/pass_diff.md 与可能的 docs/tools/index.md 导航更新;不修改任何 Python 源码。
第一步:阅读 tilelang/tools/ 中 lower_trace 与 pass_diff 的实现与 CLI,以及 docs/tools/ 现有文档结构。
本机怎么复现 / 验证:在 MacBook 上 pip install -e . 后,用 CPU 后端运行 examples/gemm 中一份简单 kernel 并调用 tilelang.tools.lower_trace / pass_diff 生成 IR dump;无需 GPU。
认领留言(英文,可直接贴到 Issue)
I'd like to add end-to-end docs for the transform-debugging tools that #3459 surfaces as missing: docs/tools/lower_trace.md and docs/tools/pass_diff.md, walking through a simple kernel (e.g. examples/gemm/example_gemm.py) on the CPU backend and showing the IR at each pass. No change to T.copy semantics — purely docs + reproducer. Before I start: is there a preferred CLI entry point (python3 -m tilelang.tools.lower_trace vs. Python API), and what's the doc-site build tool? Plan to open a PR in 2 evenings.
大致实施方案
  • 在 MacBook 上用 CPU 后端跑一个简单 kernel,调用 tilelang.tools.lower_trace(或 CLI)生成 IR 在每一步 pass 的 dump
  • 把输出贴进 docs/tools/lower_trace.md,并配一段文字解释每一步 pass 做了什么
  • 同理为 pass_diff 生成两份 IR(例如改一个 pass config 前后)并展示 diff 输出
  • 在 docs/tools/index.md 或 README 中把两份新文档加入导航
  • 本地用 python3 -m pytest testing 确认没有因 import 变化而跑挂(文档改动一般不触发,但保险起见)
可能涉及的目录或文件
  • docs/tools/ —— 新增 lower_trace.md、pass_diff.md
  • tilelang/tools/ —— 读取 lower_trace / pass_diff 的调用入口
  • examples/gemm/ 或 examples/elementwise/ —— 作为示例 kernel
验收方式
  • python3 -m pytest testing 通过(文档改动不应破坏测试)
  • 手动复现:读者按文档步骤跑同一命令能得到相似 IR dump
  • pre-commit run --all-files 通过(Markdown 需满足 lint 规则)
开工前问题与风险

向维护者确认

  • lower_trace / pass_diff 是否已有隐藏的 CLI(例如 python3 -m tilelang.tools.lower_trace),还是只能从 Python API 调用?
  • 文档站点的构建方式(MkDocs / Sphinx)与路径规范是什么?docs/tools/ 下的文件是否会自动渲染?
  • 是否需要同步更新 README 的 Tools 章节?

风险

  • 工具 CLI 在不同版本间可能接口漂移——应锁定到当前 main 的行为
  • IR dump 输出可能因环境不同而不完全一致——文档中应注明「输出可能因版本而异,以说明流程为主」
任务 4可选 · medium · Mac · 3 个晚上

[BUG] [METAL] T.copy at default coalesced_width incorrectly strides global-to-shared copies

无人认领1 条评论bug更新 2026-10-01

用到的专长:候选人在 Apple Silicon 上可直接编译、调试 Metal pipeline,能写出「无 GPU 也能跑」的 Metal 测试。

目标:在 maint/ 下新增 maint/metal/correctness_evaluation_metal.py,验证 Metal 后端能编译出正确形状的 kernel 产物(AOT 编译到 Metal shader 源码),并可选地与 CPU 后端输出做 bit-exact 或近似对照。

为什么值得长期做:#3292 是 Metal 后端 bug,且 README 2026-05-22 / 2026-07-28 已把 Metal simdgroup / cooperative-tensor GEMM 作为主线能力推出,但 maint/ 下完全没有 Metal 后端的 correctness 评估。这张卡把「Metal 后端到底对不对」做成可重复的 compile-only 回归测试,与 #3292 的修复形成闭环。

怎么介入:#3292 是 Metal bug 且无 open PR,本卡不直接修复 #3292,而是借该 bug 暴露出的「Metal 正确性验证缺失」做基础设施;engagement 建议明确说明「先补测试,后续可用它回归 #3292 类 bug」。
第一个 PR 的边界:第一个 PR 只新增 maint/metal/correctness_evaluation_metal.py 与可能的 README;不修改 src/metal 或 tilelang/metal 代码。
第一步:阅读 src/metal 与 tilelang/metal 的 codegen 入口,以及 examples/ 中 Metal 相关示例(如有)。
本机怎么复现 / 验证:在 MacBook Air 上 pip install -e . 后 python3 maint/metal/correctness_evaluation_metal.py,触发 Metal codegen 并断言产物;无需独立 GPU。
认领留言(英文,可直接贴到 Issue)
I'd like to add a Metal-backend compile-only correctness suite under maint/metal/correctness_evaluation_metal.py, mirroring the pattern of maint/gemm/correctness_evaluation.py. It will AOT-compile a small GEMM to Metal shader source on my Apple Silicon Mac and assert the emitted source has the expected shape/threadgroup keywords; optionally cross-check against the CPU backend. This does not fix #3292 directly but gives us a regression harness for it. Before I start: what is the current compile-only entry point (tilelang.tools.compile_only or Python API)? And should I cover both simdgroup and cooperative-tensor paths? Plan to open a PR in 3 evenings.
大致实施方案
  • 在 maint/metal/correctness_evaluation_metal.py 中构造一个最小 GEMM kernel,用 compile_only=True 或 equivalent 触发 Metal codegen 但不实际运行
  • 断言编译产物是合法的 Metal shader 源码(含 threadgroup、simdgroup 等关键字,形状与输入匹配)
  • 若可能,把同一 kernel 在 CPU 后端跑通,对比输出做近似对照(允许 Metal 路径只验证编译,CPU 路径验证数值)
  • 脚本头部注明需要 Apple Silicon + macOS
  • 在 MacBook 上跑通:python3 maint/metal/correctness_evaluation_metal.py
可能涉及的目录或文件
  • maint/metal/ —— 新增目录
  • tilelang/metal/ 与 src/metal/ —— 读取 codegen 入口
  • tilelang/tools/compile_only —— 可能已有的 compile-only 工具
  • examples/ 中 Metal 示例 —— 作为 kernel 写法参考
验收方式
  • python3 maint/metal/correctness_evaluation_metal.py 在 MacBook 上编译 Metal shader 成功并输出 pass
  • 编译产物含预期关键字(如 threadgroup、simdgroup)
  • 若含 CPU 对照,数值误差在合理范围
  • pre-commit run --all-files 通过
开工前问题与风险

向维护者确认

  • Metal codegen 的 compile-only 入口是什么?是否已有 tilelang.tools.compile_only 或等价 API?
  • 是否需要同时验证 simdgroup 与 cooperative-tensor 两条路径,还是先只覆盖其中一条?
  • Apple M 系列芯片的代际差异(M1 vs M5)是否会影响测试稳定性?

风险

  • Metal codegen 在 headless CI 上可能不可用——需要标注「Apple Silicon only」
  • cooperative-tensor GEMM 只支持 M5 及更新机型——需在脚本中做 capability 检测
任务 5可选 · medium · CPU · 2 个晚上

[Test][Autotune] 在 CPU 后端复现结构化 trial 报告并补全文档

无人认领0 条评论更新 2026-10-02

用到的专长:候选人熟悉训练系统性能分析,可设计可复现的调优报告流程。

目标:在 benchmark/ 下新增 benchmark/autotune_cpu/ 目录,包含一个最小 autotune 示例(如 GEMM tile 大小搜索),使用 #3393 引入的结构化 trial 报告功能输出 trials.jsonl,并在 docs 中补一段使用说明。

为什么值得长期做:#3375 已被 #3393 实现并合并(2026-10-12),但 README 与 benchmark/ 下仍缺少「如何用这个功能产出可复现调优报告」的示例。这张卡把 #3393 的功能落到 CPU 后端上,产出 benchmark/autotune_cpu/ 复现脚本与文档,验证结构化 trial 输出在 CPU 路径下可用。

怎么介入:#3375 已被 #3393 实现,本卡不重复实现功能,而是做「使用示例 + CPU 复现脚本 + 文档」;engagement 建议明确说明「#3393 已实现,本 PR 只补 benchmark 与 docs」。
第一个 PR 的边界:第一个 PR 只新增 benchmark/autotune_cpu/reproduce_structured_trial.py 与 docs/programming_guides/autotuning.md 的一段说明;不修改 tilelang/autotuner 核心代码。
第一步:阅读 #3393 的 commit(07295bba)与 tilelang/autotuner 中结构化 trial 报告的 API 入口,以及 docs/programming_guides/autotuning 现有文档。
本机怎么复现 / 验证:在 MacBook 上 pip install -e . 后 python3 benchmark/autotune_cpu/reproduce_structured_trial.py,走 LLVM CPU 后端 + 结构化 trial 报告;无需 GPU。
认领留言(英文,可直接贴到 Issue)
#3393 already landed the structured trial reporting feature. I'd like to add a CPU-backend reproducer under benchmark/autotune_cpu/reproduce_structured_trial.py that runs a small GEMM autotune with on_trial=... and writes trials.jsonl, plus a docs snippet in docs/programming_guides/autotuning.md showing how to use it. This is docs + benchmark only — no change to tilelang/autotuner core. Before I start: is the final on_trial callback signature confirmed on main, and should this be wired into CI or stay as a local reproducer? Plan to open a PR in 2 evenings.
大致实施方案
  • 在 benchmark/autotune_cpu/reproduce_structured_trial.py 中定义一个带 pass config 搜索空间的 GEMM kernel,调用 AutoTuner.run(on_trial=...) 把每条 trial 写到 trials.jsonl
  • 在 MacBook CPU 后端上跑通,确认 trials.jsonl 每条记录含 index / config / status / latency_ms / error 等字段
  • 在 docs/programming_guides/autotuning.md 中补一段「结构化 trial 报告」的使用说明,引用 benchmark/aututune_cpu/ 作为可复现入口
  • 本地运行 python3 benchmark/autotune_cpu/reproduce_structured_trial.py 并检查输出 JSONL
可能涉及的目录或文件
  • benchmark/autotune_cpu/ —— 新增目录
  • tilelang/autotuner/ —— 读取结构化 trial 报告 API
  • docs/programming_guides/autotuning.md —— 补全文档
  • #3393 的 commit 07295bba —— 确认 API 形状
验收方式
  • python3 benchmark/autotune_cpu/reproduce_structured_trial.py 在 MacBook CPU 后端上生成合法的 trials.jsonl
  • 每条 JSONL 记录含 index、config、status、latency_ms(或 error)字段
  • 文档站点(如有本地 build 命令)渲染正常
  • pre-commit run --all-files 通过
开工前问题与风险

向维护者确认

  • 结构化 trial 报告的 API 是 on_trial=lambda record: ... 还是其他形状?需确认 #3393 合并后的最终接口
  • 是否需要把 benchmark/autotune_cpu/ 接入 CI,还是作为本地复现脚本即可?
  • CPU 后端上 autotune 的搜索空间是否应限制在小规模(避免跑太久)?

风险

  • #3393 的 API 在合并后可能有微调——需以 main 分支实际代码为准
  • CPU 后端 autotune 若搜索空间太大会跑很久——应限制 trial 数量