← 所有项目

Contribution Tasks

tile-ai/tilelang

TileLang 是一个基于 TVM TIRX 的高性能 GPU/CPU 内核 DSL,正处于快速迭代的 Beta 阶段。项目将 Metal 和 CPU 视为一等公民,且核心的布局推导(Z3)和编译器 Pass(C++)完全在 Host 端执行。这与候选人资深的 GPU kernel、算子融合、端侧部署经验完美契合。候选人可以在 Mac 上完成绝大部分编译器前端、IR 降级、Metal 后端适配和性能建模的开发与验证,仅在必要时使用 Colab T4 验证 CUDA 逻辑。

当前方向:近期项目密集添加了前沿硬件(SM100/SM120、CDNA4)的特性支持,同时在完善多后端(Metal、CPU、LLVM)的对齐,并积极修复 Fuzzer 发现的编译期边界与降级问题。

★ 7446Fork 7456 个候选任务Gemini:gemini-3.1-pro-preview

更新于 2026-10-08T07:37:23+00:00 · 本次生成失败,展示上次结果 · 打开仓库 ↗

一、项目定位

TileLang (tile-lang) 是一个基于 Python 的领域特定语言(DSL),构建在 Apache TVM 编译器基础设施(TIRX)之上,专为简化高性能 GPU/CPU/加速器内核(如 GEMM、FlashAttention)的开发而设计。它的核心定位是让开发者能够使用 Pythonic 的语法进行 Tile 级别的编程,同时保留底层硬件优化的能力(如 TMA、Tensor Cores、Warp 特化)。与直接编写 CUDA 或 C++ 模板(如 CuTe)相比,它大幅提升了生产力;与同类的 Triton 相比,它不仅支持 NVIDIA 和 AMD,还拥有非常活跃的 Metal 和 CPU (LLVM) 后端。目前项目处于 Beta 阶段(Development Status :: 4 - Beta),迭代极其迅速,已经原生支持了最前沿的硬件特性(如 NVIDIA Blackwell SM100/SM120 的块缩放 MMA、AMD CDNA4 MXFP4、Apple M5 协作张量 GEMM)。 对于你(资深 GPU/训练性能工程师)而言,这个项目简直是为你量身定制的完美切入点。虽然你目前只有一台 16GB 的 Apple Silicon MacBook Air,但 TileLang 将 Metal 视为一等公民(拥有独立的 src/metal 和 tilelang/metal 目录,且最近的提交热点包含 Metal 的 32-bit atomic add 支持)。你完全可以在本地 Mac 上使用 Metal 或 CPU 后端完成算子融合、量化(INT4/FP8/FP4)或扩散模型(DiT)推理加速算子的开发、布局推导(基于 Z3)与本地测试,最后再用 Colab T4 验证 CUDA 逻辑。这里有大量你熟悉的场景(DeepSeek MLA、Block-causal attention),你的端侧部署和内核优化经验在这里可以直接转化为核心贡献。

解决什么问题、给谁用

解决的核心问题是:编写极致优化的深度学习算子(如支持动态形状的 FlashAttention、低比特量化 GEMM、复杂的 MoE 路由)通常需要手写极其冗长且难以调试的 CUDA/C++ (CuTe) 代码,且这些代码与特定硬件(NVIDIA)深度绑定,难以跨平台移植到 AMD (ROCm) 或 Apple Silicon (Metal)。

目标用户为 AI 系统架构师、GPU 内核工程师、大模型/扩散模型推理优化研究员。典型使用场景包括:为新型模型架构(如 DeepSeek V3/V4、DiT 视频生成)编写定制化的 Attention 或 GEMM 算子;在端侧设备(Mac/Windows)上实现高性能的量化推理(FP8/INT4);以及利用自动调优(Autotuning)快速探索算子的最佳 Tile 调度策略。

同类项目与差别

核心能力

能力在哪成熟度
Metal 后端代码生成与执行 (Apple Silicon 支持)src/metal, tilelang/metal, requirements-test-metal.txt成熟
CUDA 后端与前沿架构支持 (SM90/SM100/SM120, TMA, Tensor Cores)src/cuda, tilelang/cuda, examples/blockscaled_gemm_sm100成熟
CPU (LLVM) 后端执行src/cpu, tilelang/cpu成熟
ROCm 后端 (AMD RDNA3/CDNA4 支持)src/rocm, tilelang/rocm, examples/amd成熟
前沿大模型算子实现 (DeepSeek MLA/NSA/MHC, BitNet)examples/deepseek_mla, examples/deepseek_nsa, examples/bitnet-1.58b成熟
扩散模型/长文本注意力机制 (Block-causal attention)examples/block_causal_attention成熟
低比特量化矩阵乘法 (FP8, INT4, FP4)examples/gemm_fp8, examples/gemm_int4, pyproject.toml (fp4 extra)成熟
基于 Z3 的符号化布局推导与验证maint/layout_inference, tilelang/layout, pyproject.toml (z3-solver)成熟
自动调优 (Autotuning) 与 JIT 编译tilelang/autotuner, tilelang/jit成熟
WebGPU 后端src/webgpu, tilelang/webgpu需验证

阶段:Beta 阶段(根据 pyproject.toml 的 classifiers: Development Status :: 4 - Beta)。项目迭代极快,2026年7月到9月连续发布了 v0.1.12 到 v0.1.14 版本。

技术栈:Python (>=3.10);C++ (C++17/20 兼容 MSVC/Clang/NVCC);Apache TVM (apache-tvm-ffi);PyTorch (torch, torch-c-dlpack-ext);Z3 Solver (z3-solver);CMake & scikit-build-core;uv (Python 依赖管理);NVCC / NVRTC (CUDA 编译);pytest (测试框架);Sphinx (文档构建);pre-commit / Ruff / clang-format (代码规范)

规模:项目规模庞大且活跃:包含超过 1300 个文件(examples 目录 346 个,testing 378 个,src 301 个,tilelang 354 个)。提交频率极高,仅在 2026年9月16日至21日期间就有 12 次核心提交。测试覆盖广泛,testing/python 是最大的提交热点(37次)。2026年下半年保持每月发布一个 Release 的节奏(v0.1.12 到 v0.1.14)。

二、架构与代码地图

TileLang 的架构设计深受 TVM 和 Halide 的启发,但专注于 Tile 级别的编程抽象,整体可划分为五个核心层次,非常适合你这种具备底层 Kernel 编写经验但目前受限于硬件环境的工程师进行本地开发与调试: 1. 前端接入与 DSL 层 (Frontend & DSL Layer):位于 tilelang/language 和 tilelang/tileop。这是用户直接交互的层,提供高度 Pythonic 的语法(如 @T.prim_func, T.Kernel, T.gemm)。它将复杂的硬件指令(如 TMA、MMA)抽象为高级函数。你可以在 Mac 上完全无缝地编写这些 Python 代码,无需任何 GPU 环境。 2. 核心编译与 IR 变换层 (Compiler & Transform Layer):位于 src/transform 和 tilelang/transform。这是 TileLang 的“大脑”,负责将前端的 Python AST 转换为底层的 TIRX(基于 TVM 的 IR)。这里包含了大量针对高性能计算的 Pass(如 Warp 特化、流水线调度、向量化)。 3. 布局推导层 (Layout Inference Layer):位于 src/layout 和 maint/layout_inference。这是该项目最硬核的特性之一,使用 Z3 SMT 求解器(z3-solver)自动推导 Shared Memory 和 Register 的最优布局,避免 Bank Conflict。这部分逻辑完全在 CPU 上运行,是你用 Mac 贡献核心代码的绝佳切入点。 4. 多后端代码生成层 (Backend CodeGen Layer):位于 src/backend、src/cuda、src/metal 和 src/cpu。将优化后的 TIRX 翻译为目标硬件的源码(CUDA C++、Metal Shading Language 或 LLVM IR)。项目将 Metal 视为一等公民,这意味着你的 Apple Silicon Mac 是原生支持的开发目标。 5. JIT 编译与运行时层 (JIT & Runtime Layer):位于 tilelang/jit 和 tilelang/autotuner。负责调用本地编译器(如 Apple 的 metal 工具链或 Windows/Linux 的 nvrtc)将源码编译为动态链接库,并通过 TVM FFI 暴露给 PyTorch。你可以在本地完成 Metal 算子的 JIT 编译和正确性验证,最后再把代码丢到 Colab T4 上验证 CUDA 逻辑。

用户接口层编译与变换层后端生成层运行时层SOTA_Examples → Python_DSL:调用语法调用语法DL_Operators → Python_DSL:调用语法Python_DSL → IR_Transform:AST/IRAST/IRIR_Transform → Layout_Inference:查询布局IR_Transform → Metal_Backend:TIRXTIRXIR_Transform → CUDA_Backend:TIRXTIRXMetal_Backend → JIT_Compiler:MSL源码MSL源码CUDA_Backend → JIT_Compiler:CUDA源码CUDA源码Python_DSLPython_DSLDL_OperatorsDL_OperatorsSOTA_ExamplesSOTA_ExamplesIR_TransformIR_TransformLayout_InferenceLayout_InferenceMetal_BackendMetal_BackendCUDA_BackendCUDA_BackendJIT_CompilerJIT_Compiler
TileLang 从 Python DSL 定义、Z3 布局推导到多后端 JIT 编译的完整架构数据流。
模块 / 路径职责 · 入口 · 依赖
Python_DSL
tilelang/language
中
提供 Pythonic 的 Tile 级别编程接口,将用户代码解析为抽象语法树 (AST)。
入口:T.prim_func, T.Kernel, T.Parallel
依赖:tilelang/ir.py
最近提交热点之一,刚增加了对 Python iterables 和 comprehensions 的支持。
DL_Operators
tilelang/tileop
中
内置的高级深度学习算子库,封装了常用的 GEMM、Attention 等模板。
入口:需验证 (如 tilelang/tileop/gemm.py)
依赖:Python_DSL
适合作为学习 DSL 语法的官方参考实现。
IR_Transform
src/transform
大
核心编译器 Pass,负责将高级 IR 降级、优化(如 TMA 降级、Warp 特化、原子操作向量化)。
入口:src/transform/ (C++ Passes)
依赖:TVM TIRX
提交热点,近期优化了目标地址的原子向量宽度规划。C++ 编写,可在 Mac 本地编译调试。
Layout_Inference
maint/layout_inference
中
基于 Z3 求解器的内存布局自动推导系统,解决 Shared Memory 冲突。
入口:maint/layout_inference/oracle.py, maint/layout_inference/run.py
依赖:z3-solver
纯 CPU 逻辑,非常适合你在 Mac 上进行算法优化和测试。
Metal_Backend
src/metal
中
Apple Silicon 的代码生成器,将 TIRX 翻译为 Metal Shading Language。
入口:需验证 (如 src/metal/codegen_metal.cc)
依赖:IR_Transform
近期热点,刚支持了 32-bit integer atomic add。这是你在本地验证算子逻辑的核心依赖。
CUDA_Backend
src/cuda
大
NVIDIA GPU 的代码生成器,支持最前沿的 SM90/SM100/SM120 特性。
入口:需验证 (如 src/cuda/codegen_cuda.cc)
依赖:IR_Transform
包含大量前沿硬件优化(如 SM120 NVF4 block-scaled MMA),可在 Colab T4 上做基础验证。
JIT_Compiler
tilelang/jit
小
即时编译生成的源码,并将其封装为 PyTorch 可调用的函数。
入口:tilelang.compile
依赖:Metal_Backend, CUDA_Backend, apache-tvm-ffi
支持跨主机 CUDA 二进制缓存,本地 Mac 会调用 Metal 工具链。
SOTA_Examples
examples
极大
前沿模型(DeepSeek V3/V4, DiT, BitNet)的极致优化算子参考实现。
入口:examples/deepseek_mla/, examples/blockscaled_gemm_sm100/
依赖:Python_DSL, JIT_Compiler
包含大量你熟悉的场景(MLA, Block-causal attention, FP8/INT4 量化)。
目录树(按文件数)
  • testing/ 378 个文件
    .gitkeep, conftest.py, cpp, python
  • tilelang/ 354 个文件
    __init__.py, _ffi_api.py, _typing.py, analysis, autodd.py, autotuner, backend, cache, carver, contrib, cpu, cuda, dtypes.py, engine
  • examples/ 346 个文件
    amd, analyze, attention_sink, autodd, aws, bitnet-1.58b, block_causal_attention, blockscaled_gemm_sm100, blocksparse_attention, blocksparse_gemm, cast, conftest.py, convolution, deepseek_deepgemm
  • src/ 301 个文件
    backend, config.h, cpu, cuda, ir.cc, layout, metal, op, rocm, runtime, span_utils.cc, span_utils.h, support, tl_templates
  • maint/ 72 个文件
    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
  • benchmark/ 19 个文件
    blocksparse_attention, compile_speed, mamba2, matmul, matmul_fp8, matmul_metal
  • .agents/ 17 个文件
    skills
  • .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
  • .claude/ 8 个文件
    skills
  • images/ 8 个文件
    MatmulExample.png, MatmulExample.svg, 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
  • cmake/ 6 个文件
    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 中,一次典型的算子开发与执行流程(从定义到输出)如下,整个过程的设计使得你完全可以在 Apple Silicon Mac 上完成 95% 的工作: 1. 算子定义 (Frontend):用户在 Python 中使用 @T.prim_func 装饰器定义 Kernel。例如,编写一个用于 DiT 的 Block-causal attention 算子。代码中使用 T.alloc_shared 声明共享内存,使用 T.Parallel 定义线程块,使用 T.gemm 或 T.tma_copy 描述计算和数据搬运。此时,Python 代码被解析为 TileLang 的内部 AST。 2. 布局推导 (Layout Inference):这是 TileLang 的核心特色。当定义了多维 Tensor 并在 Shared Memory 中进行复杂访存时,编译器会调用基于 z3-solver 的布局推导引擎(maint/layout_inference)。它会在 CPU 上通过 SMT 求解,自动计算出最优的 Swizzle 布局以避免 Bank Conflict。这一步完全不依赖 GPU,你在 Mac 上可以极速完成推导与验证。 3. IR 降级与变换 (Transform):AST 被转换为底层 TIRX(TVM IR)。随后,src/transform 中的一系列 C++ Pass 开始工作。例如,如果你启用了 Warp 特化(TL_DISABLE_WARP_SPECIALIZED: False),Pass 会自动将 Producer(数据搬运)和 Consumer(MMA 计算)分离到不同的 Warp 中;或者进行原子操作的向量化合并(如最近提交的 Plan atomic vector widths)。 4. 后端代码生成 (CodeGen):根据 tilelang.compile 传入的 target 参数,IR 被路由到对应的后端。在你的 Mac 上,指定 target="metal",src/metal 会将 TIRX 翻译为 Metal Shading Language (MSL)。如果你在 Colab 上,src/cuda 会生成包含 PTX/CUDA C++ 的代码。 5. JIT 编译与执行 (JIT & Runtime):tilelang/jit 接管生成的源码。在 Mac 上,它调用 Apple 的 Metal 编译器生成动态库;在 CUDA 环境下,调用 nvrtc。编译后的二进制文件通过 apache-tvm-ffi 包装为 PyTorch 函数。最后,你传入 PyTorch Tensor(在 Mac 上是 MPS 或 CPU Tensor),Kernel 被分发到硬件执行,结果写回输出 Tensor。你可以利用 tools/lower_trace 观察每一步的 IR 变化,或者用 tools/pass_diff 对比优化效果。

1定义 KernelPython_DSL · tilelang/language2解析为 IRIRBuilder · tilelang/ir.py (需验证)3推导布局Layout_Inference · maint/layout_inference/oracle.py4IR 变换IR_Transform · src/transform5生成源码Metal_Backend · src/metal6JIT 编译JIT_Compiler · tilelang/jit7执行与输出TVM FFI · tilelang/_ffi_api.py

关键类型与函数

名称路径用途
T.prim_functilelang/languagePython 装饰器,用于标记和解析 TileLang Kernel 的入口函数。
T.Kerneltilelang/language上下文管理器,用于定义 Kernel 的 Grid 和 Block 维度(如 T.ceildiv(n, block_n))。
T.Tensortilelang/language定义多维张量的数据结构,支持指定形状和数据类型(如 T.float16, T.int4)。
tilelang.compiletilelang/jitJIT 编译器的核心入口,将 Python AST 编译为可执行的 PyTorch 函数,支持指定 target 和 pass_configs。
T.alloc_sharedtilelang/language在 Kernel 内部显式分配 Shared Memory(或 Metal 的 Threadgroup Memory)。
T.tma_copytilelang/language硬件加速的异步内存拷贝指令(Tensor Memory Accelerator),在 Metal/CPU 后端会自动降级为等效的标量/向量拷贝。
T.gemmtilelang/language矩阵乘法 Intrinsic,根据后端自动映射到 CUDA Tensor Cores (MMA) 或 Metal simdgroup_matrix。
PassConfigKeytilelang/transform (需验证)配置编译器 Pass 行为的枚举/字典键,如 TL_DISABLE_WARP_SPECIALIZED。
T.mbarrier_arrive / T.mbarrier_wait_paritytilelang/language异步操作(如 TMA)的底层同步屏障原语。

扩展点

最近在动的地方

建议阅读顺序

  1. README.md
  2. docs/index.md
  3. examples/quickstart.py
  4. tilelang/language/
  5. benchmark/matmul_metal/benchmark_matmul_metal.py
  6. src/metal/
  7. src/transform/
  8. maint/layout_inference/
  9. examples/deepseek_mla/

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

安装

  1. git clone --recurse-submodules https://github.com/tile-ai/tilelang.git && cd tilelang
  2. uv venv --seed .venv && source .venv/bin/activate
  3. uv pip install -r requirements-dev.txt -r requirements-test-metal.txt
  4. python3 -m pip install --no-build-isolation --verbose --editable .

哪些路径能真跑

作为只有 Apple Silicon 的开发者,你的核心优势在于 TileLang 将 Metal 和 CPU 视为一等公民。你可以直接在 Mac 上完整运行 src/metal 和 src/cpu 的代码生成与 JIT 编译。更重要的是,TileLang 最硬核的 Shared Memory/Register 布局推导(基于 Z3 求解器,位于 maint/layout_inference/ 和 src/layout/)完全在 CPU 上执行。你可以本地编写针对 DeepSeek V3/V4 或 SM100 的复杂 CUDA 算子(如 examples/deepseek_v32/ 或 examples/blockscaled_gemm_sm100/),利用 Pass Visualizer 和 IR Lower Trace 在本地观察 TIRX 的 lowering 过程和 TMA 访存优化,仅在最后一步将代码扔到 Colab T4 上验证 nvrtc 编译和执行正确性。

最小可运行

  1. python3 examples/matmul_metal/benchmark_matmul_metal.py
  2. python3 maint/layout_inference/run.py
  3. python3 examples/plot_layout/layout_swizzle.py

测试

测试框架使用 pytest,主要目录为 testing/python(近期有 37 个 commit,是核心热点)和 testing/cpp。在 Mac 上为了避免触发 CUDA/ROCm 相关的 JIT 失败,建议使用过滤命令运行 CPU 和 Metal 子集:pytest testing/python -k "not cuda and not rocm"。本地执行该子集通常在几分钟内完成。

调试

CI

CI 运行在 GitHub Actions (.github/workflows) 上。PR 提交前会被严格的 Lint 卡住,必须在本地执行 pre-commit run --all-files(包含 Ruff 检查)以及 bash format.sh(包含 clang-format 检查)。此外,CI 会运行 maint/scripts/regression_all.py 确保没有性能回退。

坑

四、维护者与社区

Release 频率非常稳定,约为每月发布一个版本(证据:v0.1.12 发布于 2026-07-08,v0.1.13 发布于 2026-08-03,v0.1.14 发布于 2026-09-02)。日常提交极其活跃,仅 2026年9月16日至21日期间就有 12 个 commits 被合并,几乎每天都有代码合入主分支。

谁角色依据
Lei Wang核心维护者 / 项目发起人在 pyproject.toml 的 maintainers 字段中被明确列出(email: leiwang1999@outlook.com)。
KellyFrog核心贡献者被指派处理 Issue #3178([RFC] Add explicit CUDA cache policies for normal global-memory copies)。
Dexterai核心贡献者被指派处理 Issue #3120(Lower T.any_of / T.all_of via warp vote intrinsics)。

流程与 Review 风格

贡献流程规范且重度依赖自动化工具。要求先在 Issue 中提问或讨论。本地开发推荐使用 uv venv 配置环境,并强制要求安装 pre-commit 钩子(包含 Ruff)。代码提交前必须通过 bash format.sh 格式化(C/C++ 使用 clang-format)。对于 C++ API 的修改,需遵循 docs/developer_guide/cpp_style.md,并运行 python3 maint/scripts/audit_cpp_api_style.py 进行审计。提交 PR 时明确要求附带测试和文档更新,本地测试通过 python3 -m pytest testing 执行。

响应速度极快,务实且极度注重编译器工程质量。从 Issue 列表看,最近几天(9月17日-21日)提交的 Issue(如 #3261, #3240)都有维护者快速跟进评论。PR 合并迅速(如 #3245, #3257 等)。社区对 Fuzzer 报出的边缘情况非常重视,Issue 列表中有大量详细追踪的 ice-on-valid-code(内部编译错误)和 wrong-code(错误代码生成)问题,体现了严谨的 Review 风格。

渠道

这里的规矩

维护者现在最想要的帮助

五、切入方案

建议长期负责:Metal 后端核心维护者 (Metal Backend Maintainer) & 布局推导架构师 (Layout Inference Architect)
这是完美契合你当前「零 NVIDIA GPU + 仅有 Apple Silicon Mac」硬约束的战略方向。Metal 后端(src/metal, tilelang/metal)允许你在本地完成 100% 的开发、编译和运行验证,将你的算子优化经验直接变现;而布局推导(src/layout, maint/layout_inference)作为 TileLang 的核心编译器大脑,完全依赖 CPU 上的 Z3 求解器运行。你可以凭借对 CuTe/TMA 的深刻理解,在 Mac 上为最顶级的 Blackwell 架构编写内存布局推导算法,从而在不触碰物理 GPU 的情况下,掌控该项目最硬核的编译器模块。

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

缺口依据为什么是你
Metal 后端对前沿复杂算子(如 Block-causal attention、DiT 算子)的支持落后于 CUDA仓库中有大量针对 CUDA 的前沿算子(如 examples/block_causal_attention/、examples/deepseek_mla/),但 Metal 相关的提交(如 [Metal] Support 32-bit integer atomic add (#3211))显示 Metal 后端仍在补齐基础能力阶段。你具备深厚的算子融合与扩散模型推理加速经验,清楚这些复杂算子的数学逻辑与访存瓶颈。你可以直接在 Mac 上将这些高级算子的逻辑降级(Lowering)并适配到 Metal 后端,实现端侧加速。
新一代硬件(SM100/SM120)的 Shared Memory/Register 布局推导规则需要大量 CPU 端的 Z3 约束编写项目近期热点包括 examples/blockscaled_gemm_sm100 和 examples/gemm_sm120,且 maint/layout_inference/ 目录负责核心的布局推导(基于 Z3 求解器)。你精通 CuTe/CUTLASS 和 TMA/MMA 机制,完全理解 Bank Conflict 和 Swizzle 逻辑。布局推导是纯 CPU 计算任务,你可以用 Mac 编写和调试 Z3 约束,为最前沿的 NVIDIA 硬件贡献核心大脑。
端侧量化(FP4/FP8/INT4)在 Metal/CPU 后端的端到端验证与性能调优项目支持多种量化格式(examples/dequantize_gemm/quantize/nvfp4.py、examples/gemm_int4),但主要集中在 Hopper/CDNA4 架构。Metal 端的低比特量化支持有待完善。你拥有丰富的端侧部署和量化经验,可以在 Mac 上利用 Metal 后端开发 INT4/FP8 的反量化 GEMM 算子,并进行本地正确性验证。
C++ 核心 API 的代码规范审计与重构存在 maint/scripts/audit_cpp_api_style.py 脚本,且 docs/developer_guide/cpp_style.md 明确了 C++ 规范,说明项目正处于从快速迭代向工程规范化过渡的阶段。作为资深架构师,你具备优秀的 C++ 工程素养。这类重构工作完全依赖静态分析和本地编译,不需要任何 GPU 即可完成,是熟悉底层代码库(src/)的绝佳途径。
Host 端 Kernel Launch 开销的跨平台(CPU/Metal)基准测试存在 maint/host_checks/cutedsl_launcher_overhead.py 等脚本用于测试 CUDA/CuTe 的 Launch 开销,但缺乏针对 Metal 或纯 CPU 调度的系统性开销分析。你熟悉 PyTorch 算子底层调度与分布式训练性能,对 Launch Overhead 极其敏感。你可以在本地扩展这些 Host Check 脚本,完善非 CUDA 平台的性能基准。
非 CUDA 平台(Metal/CPU)的自动化测试覆盖率近期 testing/python 有 37 个 commit,是绝对的热点。但大量测试可能默认依赖 CUDA 环境,导致在无 GPU 环境下容易触发 Skip 或 Fail。你可以通过 `pytest -k

第 1–30 天:看懂并露面

第 31–60 天:稳定产出

第 61–90 天:接管一块

第一批 PR

题目范围为什么安全
[Testing] Add CPU/Metal fallback tests for layout inference casesmaint/layout_inference/cases/ 和 testing/python/纯 Python 和 CPU 逻辑,不依赖任何 GPU 硬件。通过增加测试覆盖率,既能熟悉 Z3 求解器的输入输出,又能为项目提供高价值的 CI 保护,极易被 Maintainer 接受。
[Refactor] Fix C++ API style warnings reported by audit scriptsrc/ 目录下的 C++ 头文件和源文件基于官方提供的 maint/scripts/audit_cpp_api_style.py 脚本进行修复,属于静态分析驱动的重构。只需确保本地 C++ 编译通过即可,风险极低,且能展示你良好的工程素养。
[Metal] Port elementwise/gemv examples to explicitly test Metal backendexamples/elementwise/ 和 examples/gemv/这些基础算子逻辑简单,主要目的是验证 Metal 后端的连通性和正确性。你可以在 Mac 上直接运行并验证结果,无需担心复杂的硬件指令兼容性问题。
[Host] Enhance cutedsl_launcher_overhead.py to support CPU mock profilingmaint/host_checks/cutedsl_launcher_overhead.py这是一个独立的 Python 性能分析脚本。通过添加对 CPU/Metal 的 Mock 支持,你可以熟悉 TileLang 的 JIT 编译和 Host 端调用链路,且不会影响核心编译器逻辑。

怎么知道自己站住了

风险与对策
  • 缺乏 NVIDIA GPU 导致无法验证 src/cuda 目录下的 CodeGen 逻辑,容易引入回归错误 (Regression)。:严格依赖 CI 系统。在本地开发时,利用 Pass Visualizer 和 IR Lower Trace 观察 TIRX 的变化,确保逻辑正确。最终提交前,利用 Colab T4 进行快速的 Smoke Test。
  • Z3 求解器在处理复杂 Layout 约束时,可能在 Mac 上出现性能瓶颈或 Timeout。:在本地开发时,先使用小规模的 Tile Size 进行调试。如果遇到性能问题,可以利用 Python 的 cProfile 分析 maint/layout_inference/ 的性能,甚至提交优化 Z3 约束生成的 PR。
  • Metal 后端的硬件特性(如 SIMD Group)与 CUDA(如 Warp/Thread Block)存在差异,直接移植复杂算子可能失败。:深入阅读 Apple Metal Shading Language 文档,理解其与 CUDA 的映射关系。在 TileLang 中,利用多后端方言(Multi-backend language dialect,PR #2734)的特性,编写针对 Metal 的特定 Lowering Pass。
  • 在 Mac 上编译依赖 TVM FFI 的 C++ 扩展时,可能遇到环境配置或动态链接库问题。:严格遵循 pyproject.toml 中的依赖要求,确保安装了 torch-c-dlpack-ext 和 setuptools。如果遇到问题,优先查阅 docs/get_started/Installation.md 中的 macOS 源码编译指南。

六、怎么介入这个项目

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

建议顺序

成为长期维护者的路径
  • 通过修复 Fuzzer 发现的编译期 ICE 和边界检查问题,熟悉 TileLang 的 TIRX 降级流程和 C++ 核心 API。
  • 通过完善 Metal 后端适配和 MLX 集成,成为端侧部署和 Apple Silicon 支持的核心贡献者。
  • 通过优化 TensorCorePolicy 和实现 LowerMagicDiv,深入参与项目的性能建模和底层编译器优化,最终成为核心架构师。

七、任务卡

任务 1优先 · easy · CPU · 1个晚上

[BUG][Fuzzer][ice-on-valid-code] T.import_source prelude with no trailing newline fuses with the next #include, so nvcc rejects the CUDA instead of compiling

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

用到的专长:C++ 工程素养与编译器前端基础。

目标:确保 T.import_source 注入的 C 代码在生成时始终以换行符结尾,防止与后续的预处理指令粘连。

为什么值得长期做:修复基础的 Codegen 格式问题,是熟悉 TileLang C++ 代码生成器(CodeGen)入口的绝佳起点。

怎么介入:Issue 开放且无人认领,可直接留言认领并提交 PR。
第一个 PR 的边界:仅修改 Codegen 中输出 prelude 的几行代码,并添加一个 Python 测试。
第一步:定位 C++ Codegen 中处理 import_source 的逻辑,在输出字符串后强制追加换行符。
本机怎么复现 / 验证:在 Mac 上编写包含 T.import_source("__device__ int my_helper(int x) { return x + 1; }") 的 Python 脚本,调用 tilelang.compile 并打印生成的源码,观察是否粘连。
认领留言(英文,可直接贴到 Issue)
I can take this. I will locate the C++ codegen logic for import_source and ensure a newline is appended to the prelude string before emitting subsequent directives. I'll also add a Python test to verify the generated source. Expect a PR in a day or two.
大致实施方案
  • 在 src/ 目录下搜索处理 import_source 或 prelude 的 C++ 代码(可能是 src/target/ 或 src/codegen/)。
  • 在将 prelude 字符串写入输出流(stream << prelude)的地方,检查并确保追加 \n。
  • 在 testing/python/ 下新增一个测试用例(新增),使用不带换行符的 T.import_source,并断言生成的源码中包含正确的换行。
可能涉及的目录或文件
  • src/target/source/codegen_c.cc
  • src/target/source/codegen_cuda.cc
验收方式
  • 在 Mac 上运行新增的 Python 测试,通过 tilelang.compile 生成代码并检查字符串内容,确保 #if defined(_MSC_VER) 前有换行符。
开工前问题与风险

向维护者确认

    风险

    • 可能存在多个 Codegen 后端(如 Metal、ROCm)复用了该逻辑,需确保修改在基类中生效或覆盖了所有相关后端。
    任务 2优先 · easy · CPU · 1个晚上

    [BUG][Fuzzer][wrong-code] CuTeDSL codegen emits Python floor // for integer truncdiv (DivNode), giving wrong results for negative operands

    无人认领2 条评论更新 2026-08-27

    用到的专长:编译器 Codegen 与数值正确性分析。

    目标:将 CuTeDSL 后端对 DivNode 的代码生成从 Python 的 // 修改为 C 语义的 /,以保证负数截断除法的正确性。

    为什么值得长期做:修复算子数值正确性问题,深入理解 TileLang 的 AST 节点(DivNode)到目标语言的映射逻辑。

    怎么介入:Issue 开放且无人认领,可直接留言认领并提交 PR。
    第一个 PR 的边界:修改 CuTeDSL Codegen 的一个 Visitor 方法,并添加对应的字符串匹配测试。
    第一步:在 CuTeDSL 的 Codegen 中找到 VisitExpr_(const DivNode* op, ...),修改其输出格式。
    本机怎么复现 / 验证:在 Mac 上运行 Issue 中的复现脚本,指定 target="cutedsl",打印生成的代码,观察是否生成了 //。
    认领留言(英文,可直接贴到 Issue)
    I'd like to fix this. I will update the CuTeDSL codegen visitor for DivNode to emit / instead of // to match C truncation semantics. I'll also add a unit test to verify the emitted string for negative operands. PR should be ready within a day.
    大致实施方案
    • 在 src/ 目录下搜索 CuTeDSL 的代码生成器(如 codegen_cutedsl.cc)。
    • 定位到处理 DivNode 的 Visitor 方法。
    • 将输出的运算符从 // 更改为 /。
    • 在 testing/python/ 下新增测试(新增),针对负数操作数验证生成的 CuTeDSL 源码是否使用了 /。
    可能涉及的目录或文件
    • src/target/source/codegen_cutedsl.cc
    验收方式
    • 在 Mac 上运行新增的测试,指定 target="cutedsl",断言生成的代码字符串中包含 / 而不是 //。
    开工前问题与风险

    向维护者确认

      风险

      • 需确认 CuTeDSL 环境下 / 是否确实执行 C 语义的截断除法(通常是的,因为底层是 C++)。
      任务 3优先 · medium · CPU · 2-3个晚上

      [BUG][Fuzzer][accepts-invalid] Region ops (T.fill/T.clear/T.copy) accept a sliced region whose min+extent exceeds the buffer, silently emitting an out-of-bounds write

      无人认领1 条评论更新 2026-09-06

      用到的专长:编译器前端与静态分析。

      目标:在 Fill、Clear、Copy 等 Region 操作的构造函数或验证逻辑中,增加 min + extent <= shape 的静态边界检查。

      为什么值得长期做:完善编译器前端的边界检查,防止生成危险的越界内存访问代码,是提升框架鲁棒性的关键。

      怎么介入:Issue 开放且无人认领,可直接留言认领并提交 PR。
      第一个 PR 的边界:修改 Region ops 的 C++ 构造/验证逻辑,添加 Python 端的异常捕获测试。
      第一步:定位 T.fill 等操作在 C++ 端的构造函数,补充缺失的边界校验逻辑。
      本机怎么复现 / 验证:在 Mac 上运行 Issue 中的复现代码 T.fill(As[6:12], 7),观察是否静默编译通过并生成了越界的 C 代码。
      认领留言(英文,可直接贴到 Issue)
      I can work on this. I will enhance the static bounds check in the constructor of Region ops (Fill/Copy/Clear) to ensure min + extent <= shape per dimension, using arith::Analyzer if necessary for symbolic bounds. I'll add tests to verify the compiler rejects out-of-bounds slices. Expect a PR in 2-3 days.
      大致实施方案
      • 在 src/op/ 或 src/ir/ 目录下找到 Fill、Copy 等操作的定义(如 fill.cc)。
      • 找到现有的边界检查逻辑(目前只检查了 min >= 0 和 extent <= shape)。
      • 添加新的检查条件:对于每个维度,验证 min + extent <= shape,如果不满足则抛出编译期错误(如 LOG(FATAL) 或 Diagnostic)。
      • 在 testing/python/ 下新增测试(新增),使用 T.fill(As[6:12], 7) 验证是否正确抛出异常。
      可能涉及的目录或文件
      • src/op/fill.cc
      • src/op/copy.cc
      验收方式
      • 在 Mac 上运行新增的 Python 测试,使用 pytest.raises 捕获并验证编译期抛出的越界错误。
      开工前问题与风险

      向维护者确认

        风险

        • 如果 min 或 extent 是动态变量(符号),静态检查可能无法直接计算,需要使用 arith::Analyzer 进行边界推导。
        任务 4可选 · medium · CPU · 2-3个晚上

        [BUG][Fuzzer][ice-on-valid-code] Region-form tile atomic with a higher-rank source than destination crashes at compile with a raw IndexError

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

        用到的专长:编译器 IR 降级与张量维度推导。

        目标:修复 MakeIndices 在处理 atomic_add 时,当源 Region 秩大于目标 Region 秩时的越界崩溃,正确实现缺失维度的对齐(视为 size 1)。

        为什么值得长期做:修复 Tile 级别算子降级(Lowering)中的维度对齐 Bug,深入理解 TileLang 的 BufferRegion 和 Indices 映射机制。

        怎么介入:Issue 开放且无人认领,可直接留言认领并提交 PR。
        第一个 PR 的边界:修复 C++ 端的维度对齐逻辑,添加一个 Python 编译测试。
        第一步:定位 MakeIndices 或 atomic_add 的降级逻辑,修复维度遍历时的索引越界问题。
        本机怎么复现 / 验证:在 Mac 上编写 T.atomic_add(dst_2d, src_3d) 的 kernel,调用 tilelang.compile,观察是否抛出 IndexError。
        认领留言(英文,可直接贴到 Issue)
        I can take this. I will fix the MakeIndices logic for tile atomics to properly align dimensions when the source region has a higher rank than the destination, treating missing leading dimensions as size 1 as documented. I'll add a test case to prevent regressions. PR in a few days.
        大致实施方案
        • 在 src/transform/ 或 src/op/ 中搜索 MakeIndices 或处理 atomic_add lowering 的代码。
        • 分析崩溃原因:遍历源 Region 的维度时,直接用索引访问了目标 Region 的维度,导致 IndexError。
        • 修改对齐逻辑:从后向前对齐维度,或者对于目标 Region 缺失的前导维度,隐式视为 min=0, extent=1。
        • 在 testing/python/ 下新增测试(新增),传入 rank-3 source 和 rank-2 destination,验证编译通过且生成的代码正确。
        可能涉及的目录或文件
        • src/op/atomic.cc
        • src/transform/lower_tile_op.cc
        验收方式
        • 在 Mac 上运行新增的测试,确保 T.atomic_add(dst_2d, src_3d) 成功编译,并检查生成的 C/CUDA 代码中的索引计算是否正确。
        开工前问题与风险

        向维护者确认

          风险

          • 需要仔细阅读文档中关于“extents are aligned”的具体语义,确保是从右向左对齐还是从左向右对齐。
          任务 5可选 · medium · CPU · 2个晚上

          [BUG][Fuzzer][ice-on-invalid-code] T.gemm on a narrow dtype (int8/fp8/bfloat16) with any sub-atom block_K crashes nvcc instead of cleanly rejecting

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

          用到的专长:CUDA MMA 硬件指令约束与编译器前端。

          目标:在 T.gemm 的验证阶段,检查 block_K 是否小于对应 dtype 的 MMA K-atom,如果是则抛出清晰的编译期错误,而不是留给 nvcc 报错。

          为什么值得长期做:涉及底层 MMA 指令的硬件约束建模。在编译器前端拦截非法的硬件配置,是构建健壮的 GPU DSL 的核心。

          怎么介入:Issue 开放且无人认领,可直接留言认领并提交 PR。
          第一个 PR 的边界:在 C++ 端的 GEMM 验证逻辑中添加 K-atom 检查,并添加 Python 测试。
          第一步:在 src/cuda/op/gemm.cc 或类似的算子验证逻辑中,添加对 block_K 和 dtype atom 大小的校验。
          本机怎么复现 / 验证:在 Mac 上编写 int8 且 block_K=16 的 T.gemm,调用 tilelang.compile,检查生成的 CUDA 代码是否包含非法的 tl::mma_sync 模板参数。
          认领留言(英文,可直接贴到 Issue)
          I'd like to fix this. I will add a validation check in the GEMM lowering/op definition to ensure block_K is at least the size of the MMA K-atom for the given dtypes (e.g., 32 for int8/fp8, 16 for bf16). This will cleanly reject invalid configs before codegen. PR in 2 days.
          大致实施方案
          • 定位 T.gemm 的 C++ 验证逻辑(如 CheckMma 或 GetGemmInst)。
          • 根据操作数的 dtype 获取其 MMA K-atom 大小(如 int8/fp8 为 32,bf16/fp16 为 16)。
          • 检查当前的 block_K 是否小于该 atom 大小。如果是,使用 LOG(FATAL) 或 Diagnostic 抛出明确的错误信息(如 "block_K must be >= MMA K-atom")。
          • 在 testing/python/ 下新增测试(新增),验证传入非法的 block_K 时是否抛出预期的 Python 异常。
          可能涉及的目录或文件
          • src/cuda/op/gemm.cc
          • src/op/gemm.cc
          验收方式
          • 在 Mac 上运行新增的测试,使用 pytest.raises 捕获并验证 TileLang 抛出的编译期错误,而不是生成非法的 CUDA 代码。
          开工前问题与风险

          向维护者确认

            风险

            • 需要确保不同架构(如 SM75, SM80, SM90)下的 K-atom 大小被正确获取,避免误杀合法的配置。
            任务 6可选 · hard · CPU · 3-4个晚上

            [BUG][Fuzzer][ice-on-valid-code] 8-bit-operand sparse T.gemm_sp with K > block_K crashes ThreadSync (Cannot match type int32 vs handle) instead of compiling

            无人认领1 条评论更新 2026-09-08

            用到的专长:编译器 Pass 开发与 TVM TIRX 类型系统。

            目标:修复 ThreadSync pass 在处理 8-bit 稀疏 GEMM 且 K 维度有多次迭代时,由于类型不匹配(int32 vs handle)导致的编译期崩溃。

            为什么值得长期做:修复核心编译器 Pass(ThreadSync)中的类型推导崩溃,深入理解 TIRX 的类型系统和同步机制。

            怎么介入:Issue 开放且无人认领,可直接留言认领并提交 PR。
            第一个 PR 的边界:修复 ThreadSync Pass 中的类型匹配逻辑,添加一个 Python 编译测试。
            第一步:定位 src/transform/thread_sync.cc 中的 BinaryOpMatchTypes 或相关类型推导逻辑,找出 handle 类型的来源。
            本机怎么复现 / 验证:在 Mac 上编写一个 in_dtype=int8 且 K=256, block_K=128 的 T.gemm_sp kernel,调用 tilelang.compile,观察是否在 ThreadSync pass 抛出异常。
            认领留言(英文,可直接贴到 Issue)
            I can investigate this ICE. I will trace the ThreadSync pass to see where the int32 vs handle mismatch occurs for multi-iteration sparse GEMMs, likely fixing a missing cast or incorrect buffer offset type. I'll add a regression test for K > block_K with int8 operands. PR in a few days.
            大致实施方案
            • 在 src/transform/ 中找到 ThreadSync pass 的实现。
            • 分析崩溃日志:当 K > block_K 时,会生成循环,循环内的同步逻辑可能错误地将一个 buffer handle 与 int32 进行了二元操作(可能是偏移量计算)。
            • 修复类型匹配逻辑:在构建同步相关的 IR 节点时,确保操作数的类型一致(例如插入必要的 CastNode 或修正 handle 的使用)。
            • 在 testing/python/ 下新增测试(新增),使用 in_dtype=int8 且 K=256, block_K=128 的 T.gemm_sp,验证编译通过。
            可能涉及的目录或文件
            • src/transform/thread_sync.cc
            • src/transform/common/...
            验收方式
            • 在 Mac 上运行新增的测试,确保 tilelang.compile 成功生成代码而不抛出 InternalError。
            开工前问题与风险

            向维护者确认

              风险

              • ThreadSync 是核心 Pass,修改类型推导可能会影响其他算子,需要跑全量本地测试 pytest testing。