How ye-invisible/tilelang works
TileLang-Mesh is a Python tile-programming language and compiler/runtime extension for mesh-aware, distributed-memory accelerators, with its documented release focused on SunMMIO/SUVM. The main workflow is user-authored tile code through JIT compilation and target code generation to a compiled kernel that runs through the SunMMIO adapter. Distinct capabilities include mesh placement and communication, accelerator layouts, operation lowering, and SunMMIO pipeline scheduling. The README documents the NPU-IR toolchain integration; sampled excerpts do not show every runtime call path, so uncertain wiring is omitted.
Main entry points: lower.py Compiler lowering · kernel.py JIT kernel entry · ir.cc TileLang IR · layout_inference.cc IR transforms
Components 20
Programming frontend
- Tile language
tilelang/language/__init__.py - Mesh placement
tilelang/language/mesh_tensor.py - Communication ops
src/op/comm.cc
Compiler and layouts
- TileLang IR
src/ir.cc - Layout system
src/layout/layout.cc - Tile views
src/tileview/tileview.cc - Compiler lowering
tilelang/engine/lower.py - IR transforms
src/transform/layout_inference.cc
SunMMIO backend
- Pipeline planning
src/transform/sunmmio_pipeline_planning.cc - Resource-aware planner
src/transform/sunmmio_pipeline_planning_ilp.cc - SunMMIO codegen
src/target/sunmmio/codegen_sunmmio.cc - SunMMIO MLIR
src/target/sunmmio/sunmmio_mlir_builder.cc
JIT and runtime
- JIT kernel entry
tilelang/jit/kernel.py - SunMMIO adapter
tilelang/jit/adapter/sunmmio/adapter.py - Kernel cache
tilelang/jit/adapter/sunmmio/kernel_cache.py - Execution backend
tilelang/jit/execution_backend.py
Kernel optimization
- Autotuner
tilelang/autotuner/tuner.py - Kernel templates
tilelang/carver/template/matmul.py
Other components
- Kernel author
- NPU-IR tools
Connections 20
- Kernel author → Tile language: writes kernels
- Tile language → TileLang IR: defines operations
- Mesh placement → TileLang IR: expresses placement
- Communication ops → TileLang IR: adds communication
- TileLang IR → Compiler lowering: lowers from
- Layout system → IR transforms: informs layout inference
- Tile views → IR transforms: supports tile analysis
- Compiler lowering → IR transforms: runs passes
- IR transforms → Pipeline planning: prepares loops
- Pipeline planning → Resource-aware planner: schedules with ILP
- Compiler lowering → SunMMIO codegen: generates target code
- SunMMIO codegen → SunMMIO MLIR: builds MLIR
- SunMMIO MLIR → NPU-IR tools: uses toolchain
- Kernel author → JIT kernel entry: invokes kernel
- JIT kernel entry → Compiler lowering: compiles kernel
- JIT kernel entry → SunMMIO adapter: selects runtime adapter
- SunMMIO adapter → Kernel cache: reads and writes
- SunMMIO adapter → Execution backend: launches kernel
- Autotuner → Kernel templates: selects candidates
- Autotuner → JIT kernel entry: evaluates candidates