Can TileLang compile a kernel without a GPU? How does compile_only work?

TileLang can compile without a GPU present: run python -I -m tilelang.tools.compile_only --target cuda --output_file out.cu example.py and it writes the generated CUDA source without executing anything. The target must be a concrete backend such as cuda or rocm, since auto is rejected, and a failed compile exits with code 1 so CI can gate on it.

How do I estimate the FLOPs and memory traffic of a TileLang kernel?

TileLang's performance entry point is tilelang.tools.Analyzer.analysis(tir, device), which statically estimates FLOPs, global memory traffic and roofline time without running the kernel, so you can tell whether you are compute-bound or bandwidth-bound before tuning. Pass the target device as device, and it works on both Ascend and NVIDIA.

What should I watch when debugging TileLang compilation, and what are autodd and TL_LOWER_TRACE?

TileLang debugging works in two layers: tilelang.tools.plot_layout draws the layout view and python -m tilelang.autodd reports automatic-parallelization decisions, while TL_LOWER_TRACE turns on per-step TIR lowering logs and TILELANG_PASS_DIFF shows the difference a pass makes. The first two explain why a schedule looks wrong; the last two trace the compile pipeline.

Where does TileLang keep its compile cache, and how do I make recompiles fast?

TileLang caches compiled artifacts in ~/.tilelang/cache by default, so compiling the same kernel again reuses them and skips the rebuild. When a code change still hits a stale artifact, delete that directory to force a rebuild; clearing it after a version upgrade also prevents old and new artifacts from mixing.

Which environment variables matter when building TileLang, and how do USE_CUDA and ASCEND_HOME_PATH differ?

TileLang build-time variables split by route: on NVIDIA, USE_CUDA enables the CUDA backend, USE_ROCM and USE_METAL cover ROCm and Metal, TVM_ROOT points at the TVM source, and WITH_PIP_CUDA_TOOLCHAIN plus NO_VERSION_LABEL control toolchain source and version labels; on Ascend, CANN's set_env.sh sets ASCEND_HOME_PATH so the compiler can find Ascend libraries.

TileLang command cheatsheet: compile, analyze, debug

TileLang commands fall into four families: compile and build, performance analysis, pass debugging, and build-time environment variables. Two matter most: python -I -m tilelang.tools.compile_only --target cuda --output_file out.cu example.py, which compiles a kernel without a GPU, and tilelang.tools.Analyzer.analysis(tir, device), which estimates cost statically. Compiled artifacts live in ~/.tilelang/cache.

This is a reference page, organized as compile and build → performance and debug → environment variables and cache, with the command, its purpose and its expected output for each entry. Getting the environment set up is in Install TileLang.

What are TileLang's compile and build commands?

TileLang's compile commands split in two: compile_only, which turns a kernel script into backend source (python -I -m tilelang.tools.compile_only --target cuda --output_file out.cu example.py), and the pip or script commands that build TileLang itself from source (source).

Compile a kernel offline → compile_only. Run python -I -m tilelang.tools.compile_only --target cuda --output_file out.cu example.py. Expect generated CUDA source in out.cu; -I isolates environment variables, and --target must be a concrete backend such as cuda or rocm because auto is rejected (source).
Make failures stop a pipeline. The same command exits with code 1 when compilation fails. Expect a script or CI job to gate on that exit code, so the check drops straight into automation.
Build TileLang itself (NVIDIA). In the source tree set USE_CUDA=ON as needed, then produce a wheel with pip or python -m build -w. Expect a wheel in dist/ whose build switches match your toolchain (source).
Build TileLang itself (Ascend). Run ./build_wheel_ascend.sh [--enable-llvm], or bash install_ascend.sh to install in place. Expect a wheel in dist/ in the first case and a completed in-place install in the second (source).

How do I use TileLang's performance and debug commands?

TileLang offers Analyzer for cost estimates without running a kernel, plus several development-time debug entries: plot_layout for layout visualization, python -m tilelang.autodd for automatic parallelization, TL_LOWER_TRACE for lowering traces and TILELANG_PASS_DIFF for pass diffs (source).

Estimate FLOPs and traffic → Analyzer.analysis. Call tilelang.tools.Analyzer.analysis(tir, device). Expect FLOPs, global memory traffic and roofline time for that kernel, enough to tell whether it is compute- or bandwidth-bound (source).
See layouts and timing → plot_layout. Call tilelang.tools.plot_layout. Expect a view of the tiling and memory hierarchy that confirms blocks land where you intended.
Tune automatic parallelization → python -m tilelang.autodd. Run the module. Expect printed reasoning behind automatic parallelization decisions, useful when a schedule is not what you expected (source).
Watch the whole lowering → TL_LOWER_TRACE. Set that variable and recompile. Expect step-by-step TIR lowering logs; to diff a single pass, use TILELANG_PASS_DIFF instead.

These are development-time tools: Analyzer sets direction before tuning, while plot_layout, autodd and TL_LOWER_TRACE locate compile-time problems. None of them belongs in a production inference path.

TileLang build-time environment variables and the compile cache

TileLang's build-time variables split by route: on NVIDIA, USE_CUDA enables the CUDA backend, USE_ROCM and USE_METAL cover the others, USE_LLVM and TVM_ROOT set the compiler stack and TVM source, and WITH_PIP_CUDA_TOOLCHAIN plus NO_VERSION_LABEL control toolchain source and version labels; on Ascend, ASCEND_HOME_PATH locates Ascend libraries. Artifacts cache in ~/.tilelang/cache (source).

Variable / path	Purpose	Route
`USE_CUDA`	Enable the CUDA backend	NVIDIA
`USE_ROCM`	Enable the ROCm backend	AMD
`USE_METAL`	Enable the Metal backend	Apple
`USE_LLVM`	Enable the LLVM backend	General / Ascend builds
`TVM_ROOT`	Point at the TVM source	Source builds
`WITH_PIP_CUDA_TOOLCHAIN`	Use the pip-provided CUDA toolchain	NVIDIA
`NO_VERSION_LABEL`	Disable the wheel version label	Packaging
`ASCEND_HOME_PATH`	Point at the Ascend toolkit root	Ascend
`~/.tilelang/cache`	Compiled artifact cache	Both routes

Rebuild after changing settings. Build-time variables are read at compile time. Expect nothing to change until you rebuild; exporting alone leaves the installed artifacts as they were.
Delete the cache directory, not just files. Run rm -rf ~/.tilelang/cache. Expect the next compile to regenerate artifacts; this is the first step when a stale result survives a code change or upgrade.
Let CANN supply the Ascend variable. ASCEND_HOME_PATH is normally set by set_env.sh. Expect echo $ASCEND_HOME_PATH to return a value, without which the compiler cannot find Ascend libraries.