# ye-invisible/tilelang architecture

> Architecture diagram of the GitHub repository ye-invisible/tilelang, generated by GitDiagram from its file tree, README and a sample of source files (last updated 2026-09-29). It is an AI-made overview: check details against the source.

- Interactive diagram: https://gitdiagram.com/ye-invisible/tilelang
- Repository: https://github.com/ye-invisible/tilelang

## Overview

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: [`tilelang/engine/lower.py`](https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/tilelang/engine/lower.py) (Compiler lowering), [`tilelang/jit/kernel.py`](https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/tilelang/jit/kernel.py) (JIT kernel entry), [`src/ir.cc`](https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/src/ir.cc) (TileLang IR), [`src/transform/layout_inference.cc`](https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/src/transform/layout_inference.cc) (IR transforms)

## Diagram (Mermaid)

```mermaid
flowchart TD

subgraph group_frontend["Programming frontend"]
  node_language["Tile language<br/>[__init__.py]"]
  node_mesh["Mesh placement<br/>[mesh_tensor.py]"]
  node_communication["Communication ops<br/>[comm.cc]"]
end

subgraph group_compiler["Compiler and layouts"]
  node_ir["TileLang IR<br/>[ir.cc]"]
  node_layouts["Layout system<br/>[layout.cc]"]
  node_tileview["Tile views<br/>[tileview.cc]"]
  node_lowering["Compiler lowering<br/>[lower.py]"]
  node_transforms["IR transforms"]
end

subgraph group_sunmmio["SunMMIO backend"]
  node_sun_pipeline["Pipeline planning"]
  node_ilp_planner["Resource-aware planner"]
  node_sun_codegen["SunMMIO codegen<br/>[codegen_sunmmio.cc]"]
  node_sun_mlir["SunMMIO MLIR"]
end

subgraph group_runtime["JIT and runtime"]
  node_jit["JIT kernel entry<br/>[kernel.py]"]
  node_adapter["SunMMIO adapter<br/>[adapter.py]"]
  node_cache[("Kernel cache<br/>[kernel_cache.py]")]
  node_execution["Execution backend"]
end

subgraph group_tuning["Kernel optimization"]
  node_autotuner["Autotuner<br/>[tuner.py]"]
  node_templates["Kernel templates<br/>[matmul.py]"]
end

node_user(("Kernel author"))
node_npu_tools["NPU-IR tools"]

node_user -->|"writes kernels"| node_language
node_language -->|"defines operations"| node_ir
node_mesh -->|"expresses placement"| node_ir
node_communication -->|"adds communication"| node_ir
node_ir -->|"lowers from"| node_lowering
node_layouts -->|"informs layout inference"| node_transforms
node_tileview -->|"supports tile analysis"| node_transforms
node_lowering -->|"runs passes"| node_transforms
node_transforms -->|"prepares loops"| node_sun_pipeline
node_sun_pipeline -.->|"schedules with ILP"| node_ilp_planner
node_lowering -->|"generates target code"| node_sun_codegen
node_sun_codegen -->|"builds MLIR"| node_sun_mlir
node_sun_mlir -.->|"uses toolchain"| node_npu_tools
node_user -->|"invokes kernel"| node_jit
node_jit -->|"compiles kernel"| node_lowering
node_jit -->|"selects runtime adapter"| node_adapter
node_adapter -->|"reads and writes"| node_cache
node_adapter -->|"launches kernel"| node_execution
node_autotuner -->|"selects candidates"| node_templates
node_autotuner -->|"evaluates candidates"| node_jit

click node_language "https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/tilelang/language/__init__.py"
click node_mesh "https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/tilelang/language/mesh_tensor.py"
click node_communication "https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/src/op/comm.cc"
click node_ir "https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/src/ir.cc"
click node_layouts "https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/src/layout/layout.cc"
click node_tileview "https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/src/tileview/tileview.cc"
click node_lowering "https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/tilelang/engine/lower.py"
click node_transforms "https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/src/transform/layout_inference.cc"
click node_sun_pipeline "https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/src/transform/sunmmio_pipeline_planning.cc"
click node_ilp_planner "https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/src/transform/sunmmio_pipeline_planning_ilp.cc"
click node_sun_codegen "https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/src/target/sunmmio/codegen_sunmmio.cc"
click node_sun_mlir "https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/src/target/sunmmio/sunmmio_mlir_builder.cc"
click node_jit "https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/tilelang/jit/kernel.py"
click node_adapter "https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/tilelang/jit/adapter/sunmmio/adapter.py"
click node_cache "https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/tilelang/jit/adapter/sunmmio/kernel_cache.py"
click node_execution "https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/tilelang/jit/execution_backend.py"
click node_autotuner "https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/tilelang/autotuner/tuner.py"
click node_templates "https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/tilelang/carver/template/matmul.py"

classDef toneNeutral fill:#f8fafc,stroke:#334155,stroke-width:1.5px,color:#0f172a
classDef toneBlue fill:#dbeafe,stroke:#2563eb,stroke-width:1.5px,color:#172554
classDef toneAmber fill:#fef3c7,stroke:#d97706,stroke-width:1.5px,color:#78350f
classDef toneMint fill:#dcfce7,stroke:#16a34a,stroke-width:1.5px,color:#14532d
classDef toneRose fill:#ffe4e6,stroke:#e11d48,stroke-width:1.5px,color:#881337
classDef toneIndigo fill:#e0e7ff,stroke:#4f46e5,stroke-width:1.5px,color:#312e81
classDef toneTeal fill:#ccfbf1,stroke:#0f766e,stroke-width:1.5px,color:#134e4a
class node_language,node_mesh,node_communication toneBlue
class node_ir,node_layouts,node_tileview,node_lowering,node_transforms toneAmber
class node_sun_pipeline,node_ilp_planner,node_sun_codegen,node_sun_mlir toneMint
class node_jit,node_adapter,node_cache,node_execution toneRose
class node_autotuner,node_templates,node_user,node_npu_tools toneIndigo
```

## Components

### Programming frontend

- [Tile language](https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/tilelang/language/__init__.py): `tilelang/language/__init__.py`
- [Mesh placement](https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/tilelang/language/mesh_tensor.py): `tilelang/language/mesh_tensor.py`
- [Communication ops](https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/src/op/comm.cc): `src/op/comm.cc`

### Compiler and layouts

- [TileLang IR](https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/src/ir.cc): `src/ir.cc`
- [Layout system](https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/src/layout/layout.cc): `src/layout/layout.cc`
- [Tile views](https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/src/tileview/tileview.cc): `src/tileview/tileview.cc`
- [Compiler lowering](https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/tilelang/engine/lower.py): `tilelang/engine/lower.py`
- [IR transforms](https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/src/transform/layout_inference.cc): `src/transform/layout_inference.cc`

### SunMMIO backend

- [Pipeline planning](https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/src/transform/sunmmio_pipeline_planning.cc): `src/transform/sunmmio_pipeline_planning.cc`
- [Resource-aware planner](https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/src/transform/sunmmio_pipeline_planning_ilp.cc): `src/transform/sunmmio_pipeline_planning_ilp.cc`
- [SunMMIO codegen](https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/src/target/sunmmio/codegen_sunmmio.cc): `src/target/sunmmio/codegen_sunmmio.cc`
- [SunMMIO MLIR](https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/src/target/sunmmio/sunmmio_mlir_builder.cc): `src/target/sunmmio/sunmmio_mlir_builder.cc`

### JIT and runtime

- [JIT kernel entry](https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/tilelang/jit/kernel.py): `tilelang/jit/kernel.py`
- [SunMMIO adapter](https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/tilelang/jit/adapter/sunmmio/adapter.py): `tilelang/jit/adapter/sunmmio/adapter.py`
- [Kernel cache](https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/tilelang/jit/adapter/sunmmio/kernel_cache.py): `tilelang/jit/adapter/sunmmio/kernel_cache.py`
- [Execution backend](https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/tilelang/jit/execution_backend.py): `tilelang/jit/execution_backend.py`

### Kernel optimization

- [Autotuner](https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/tilelang/autotuner/tuner.py): `tilelang/autotuner/tuner.py`
- [Kernel templates](https://github.com/ye-invisible/tilelang/blob/tilelang_mesh_main/tilelang/carver/template/matmul.py): `tilelang/carver/template/matmul.py`

### Other components

- **Kernel author**
- **NPU-IR tools**

## Connections

- 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