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 发现的编译期边界与降级问题。
★ 7446 Fork 745 6 个候选任务 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 调度策略。
同类项目与差别 OpenAI Triton :Triton 是目前最流行的 Pythonic GPU DSL,基于自家的 MLIR 方言。TileLang 语法类似,但底层基于 Apache TVM (TIRX),且 TileLang 对多后端的支持更加显式和广泛(特别是对 Metal 和 CPU 的原生支持,而 Triton 极度偏向 CUDA/ROCm),此外 TileLang 引入了 Z3 求解器进行严格的布局和算术推导。Apache TVM (TensorIR) :TileLang 并非 TVM 的竞品,而是构建在 TVM 之上的高层抽象。TVM 的 TIR 相对底层,而 TileLang 提供了更符合直觉的 Tile 级别编程接口,最终 Lower 到 TVM 的 TIRX 表示。NVIDIA CUTLASS / CuTe :CuTe 是纯 C++ 模板库,学习曲线极陡峭。TileLang 实际上包含一个 "CuTe DSL backend",可以将 Pythonic 的 TileLang 代码编译生成 CuTe 代码,相当于 CuTe 的高阶 Python 前端。
核心能力 能力 在哪 成熟度 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/IR AST/IR IR_Transform → Layout_Inference:查询布局 IR_Transform → Metal_Backend:TIRX TIRX IR_Transform → CUDA_Backend:TIRX TIRX Metal_Backend → JIT_Compiler:MSL源码 MSL源码 CUDA_Backend → JIT_Compiler:CUDA源码 CUDA源码 Python_DSL Python_DSL DL_Operators DL_Operators SOTA_Examples SOTA_Examples IR_Transform IR_Transform Layout_Inference Layout_Inference Metal_Backend Metal_Backend CUDA_Backend CUDA_Backend JIT_Compiler JIT_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 定义 Kernel Python_DSL · tilelang/language 2 解析为 IR IRBuilder · tilelang/ir.py (需验证) 3 推导布局 Layout_Inference · maint/layout_inference/oracle.py 4 IR 变换 IR_Transform · src/transform 5 生成源码 Metal_Backend · src/metal 6 JIT 编译 JIT_Compiler · tilelang/jit 7 执行与输出 TVM FFI · tilelang/_ffi_api.py
关键类型与函数 名称 路径 用途 T.prim_func tilelang/language Python 装饰器,用于标记和解析 TileLang Kernel 的入口函数。 T.Kernel tilelang/language 上下文管理器,用于定义 Kernel 的 Grid 和 Block 维度(如 T.ceildiv(n, block_n))。 T.Tensor tilelang/language 定义多维张量的数据结构,支持指定形状和数据类型(如 T.float16, T.int4)。 tilelang.compile tilelang/jit JIT 编译器的核心入口,将 Python AST 编译为可执行的 PyTorch 函数,支持指定 target 和 pass_configs。 T.alloc_shared tilelang/language 在 Kernel 内部显式分配 Shared Memory(或 Metal 的 Threadgroup Memory)。 T.tma_copy tilelang/language 硬件加速的异步内存拷贝指令(Tensor Memory Accelerator),在 Metal/CPU 后端会自动降级为等效的标量/向量拷贝。 T.gemm tilelang/language 矩阵乘法 Intrinsic,根据后端自动映射到 CUDA Tensor Cores (MMA) 或 Metal simdgroup_matrix。 PassConfigKey tilelang/transform (需验证) 配置编译器 Pass 行为的枚举/字典键,如 TL_DISABLE_WARP_SPECIALIZED。 T.mbarrier_arrive / T.mbarrier_wait_parity tilelang/language 异步操作(如 TMA)的底层同步屏障原语。
扩展点 新模型算子 (New Kernel) :在 examples/ 目录下创建新目录(如 examples/dit_attention),使用 T.prim_func 编写算子,并通过 tilelang.compile(target="metal") 在本地 Mac 验证逻辑,最后用 PyTorch 编写参考实现进行 torch.testing.assert_close 对比。新硬件后端 (New Backend) :在 src/backend 注册新后端,并在 src/ 下新建目录(如 src/vulkan),继承并实现 TVM 的 CodeGen 接口,将 TIRX 映射为目标硬件源码。新编译优化策略 (New Transform Pass) :在 src/transform 中使用 C++ 编写新的 TVM TIR Pass(如针对特定量化格式的访存优化),通过 TVM_REGISTER_GLOBAL 暴露给 Python,并在 tilelang.compile 的 pass_configs 中启用。新布局推导规则 (New Layout Rule) :在 maint/layout_inference/ 中修改 Z3 约束模型,添加针对新硬件(如 AMD CDNA4 或 Apple M5)的 Bank 冲突避免规则。
最近在动的地方 testing/python :提交频率最高(37次)。项目处于快速迭代的 Beta 阶段,大量新特性(如 CPU dialect、Metal atomic)需要密集的单元测试覆盖。tilelang/cuda & src/cuda :核心竞争力所在。近期密集添加了 NVIDIA Blackwell (SM100/SM120) 的前沿特性,如 block-scaled GEMM 和 256-bit global load/store。tilelang/language :DSL 语法持续丰富,近期增加了对 Python iterables 和 comprehensions 的支持,提升前端编程体验。src/transform :编译器核心优化层。近期热点包括“规划目标地址的原子向量宽度”等底层访存优化 Pass。examples/kpool :前沿模型跟进极快,近期连续提交了 GLM-5.3 k-pool 相关的 Top-K transform 和 FP8 优化算子。src/metal :Apple Silicon 支持正在积极完善,近期刚合并了 32-bit integer atomic add,证明 Metal 是一等公民,非常适合 Mac 开发者介入。
建议阅读顺序 README.md docs/index.md examples/quickstart.py tilelang/language/ benchmark/matmul_metal/benchmark_matmul_metal.py src/metal/ src/transform/ maint/layout_inference/ examples/deepseek_mla/ 三、本地跑起来(没有 GPU 的 Mac) 安装 git clone --recurse-submodules https://github.com/tile-ai/tilelang.git && cd tilelanguv venv --seed .venv && source .venv/bin/activateuv pip install -r requirements-dev.txt -r requirements-test-metal.txtpython3 -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 编译和执行正确性。
最小可运行 python3 examples/matmul_metal/benchmark_matmul_metal.pypython3 maint/layout_inference/run.pypython3 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"。本地执行该子集通常在几分钟内完成。
调试 环境变量:使用 TILELANG_DISABLE_CACHE=1 强制重新编译,避免命中跨主机的 CUDA/Metal 二进制缓存。 IR 调试:在 tilelang.compile 中传入 pass_configs(如 TL_DISABLE_WARP_SPECIALIZED: True),结合 README 中提到的 Pass Visualizer 和 IR Lower Trace,可以在 Mac 上无卡观察 AST 到 TIRX 的每一步 Pass 变换。 C++ 规范检查:修改底层 IR 或 Metal CodeGen 后,运行 python3 maint/scripts/audit_cpp_api_style.py 检查 API 命名和头文件边界。 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 确保没有性能回退。
坑 Z3 版本地雷:pyproject.toml 严格限制了 z3-solver>=4.13.0,<4.15.5,千万不要为了解决本地环境冲突随意升级 Z3,否则会导致核心的布局推导(Layout Inference)逻辑直接崩溃。 Mac PyTorch 限制:在 Apple Silicon 上,pyproject.toml 强制要求 torch>=2.4,如果你的虚拟环境 PyTorch 版本过低会导致安装失败或 Metal JIT 报错。 CUDA 算子本地执行陷阱:虽然你可以在 Mac 上编写 T.tma_copy 或 T.mma_gemm_blockscaled 并成功生成 TIRX,但如果尝试调用执行,会因为缺少 nvrtc 和真实硬件而崩溃。请务必将这类算子的验证剥离到 Colab T4 上进行。 四、维护者与社区 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 风格。
渠道 https://discord.gg/TUrHyJnKPG 这里的规矩 提 Issue 前必须先搜索是否已有类似问题,并提供包含最小复现步骤的 Code snippet 和清晰的解释。 提交 PR 时,如果适用,必须包含 tests 和 docs(Please include tests and docs with every pull request if applicable!)。 修改 C++ API 时,需仔细检查 FFI 可见性、生成的代码和外部 API 约束,绝不能盲目使用 audit 脚本进行全局重命名(The audit is a review aid, not a blanket rename command)。 维护者现在最想要的帮助 完善 Metal 后端生态:例如 Issue #3222 提出的 MLX adapter for Metal kernels (vllm-metal),这简直是为你这台 Mac 量身定制的完美高价值切入点。 修复编译器后端的 Edge Cases:社区目前积压了大量由 Fuzzer 发现的 Bug(如 #3034, #2625, #2979 等 ice-on-invalid-code 或 wrong-code),这些纯编译期的 AST/IR 转换 Bug 修复完全可以在你的 CPU 环境下完成。 基础指令与类型优化:如 Issue #3120(Lower T.any_of / T.all_of via warp vote intrinsics,带有 good first issue 标签)和 Issue #859(Float32 数据类型修复)。 新硬件后端支持:如 Issue #1293 提到的 Hexagon Backend with HMX Supports(带有 help wanted 标签)。 七、任务卡 任务 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 Issue #3020 · [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 Issue #2973 · [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 Issue #2979 · [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 Issue #3023 · [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 Issue #3032 · [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 Issue #3037 · [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。