Compile-Only Tool¶

python -m tilelang.tools.compile_only lowers a TileLang program to inspectable kernel source without running it. No GPU, CUDA driver, or device compilation is required: the tool stops after tilelang.lower(..., enable_device_compile=False) and writes the generated kernel source to a file. It backs the TileLang integration in Compiler Explorer and is equally useful for inspecting locally what lower() produces for a kernel.

Requirements¶

The tool ships with the standard TileLang installation and has no optional dependencies. The default c target works on any machine, including CPU-only hosts. --target cuda requires a wheel whose CUDA codegen FFI is compiled in: Linux wheels include it and can emit CUDA source without a GPU present, while macOS Metal wheels do not and fail with a clear CUDA codegen FFI missing message instead of an AttributeError.

Quick Start¶

Write a self-contained example that defines a kernel but never launches it:

# example.py
import tilelang
import tilelang.language as T


@tilelang.jit
def add(A, B):
    N = 64
    A: T.Tensor((N,), T.float16)
    B: T.Tensor((N,), T.float16)
    C = T.empty((N,), T.float16)
    with T.Kernel(1, threads=64):
        for i in T.Parallel(N):
            C[i] = A[i] + B[i]
    return C

Compile it to CPU C source:

python -I -m tilelang.tools.compile_only --output_file out.c example.py

out.c then contains the generated kernel source:

// tilelang target: {"kind":"c","tag":"","keys":["cpu"],"host":{"kind":"c","tag":"","keys":["cpu"]}}
#include <tl_templates/cpp/common.h>

#ifdef __cplusplus
extern "C"
#endif
int32_t add_kernel(half* A, half* B, half* C);
#ifdef __cplusplus
extern "C"
#endif
int32_t add_kernel(half* A, half* B, half* C) {
  for (int32_t i = 0; i < 8; ++i) {
    *(half8*)(C + (i * 8)) = (*(half8*)(A + (i * 8)) + *(half8*)(B + (i * 8)));
  }
  return 0;
}

CUDA source generation works the same way and still needs no GPU; the architecture is pinned instead of detected:

python -I -m tilelang.tools.compile_only --target cuda --output_file out.cu example.py

extern "C" __global__ void __launch_bounds__(64, 1) add_kernel(const half_t* __restrict__ A, const half_t* __restrict__ B, half_t* __restrict__ C) {
  C[((int)threadIdx.x)] = (A[((int)threadIdx.x)] + B[((int)threadIdx.x)]);
}

How the Input Is Processed¶

The input file is imported as a regular Python module, so its top-level code runs. Keep examples compile-only: define kernels, do not launch them or allocate tensors at import time. Running the CLI with python -I (isolated mode) is recommended, and is how Compiler Explorer invokes it.
The first @tilelang.jit function or PrimFunc in module order is selected. Objects are inspected in definition order, so a PrimFunc defined before a @tilelang.jit kernel wins.
The kernel is lowered but never device-compiled or executed, so no tensor arguments are needed and no GPU is touched.

CLI Reference¶

python -I -m tilelang.tools.compile_only [--target TARGET] --output_file OUTPUT input_file

Argument	Meaning
`input_file`	Compile-only example file. Imported, never launched.
`--output_file`	Required destination for the generated kernel source. Must not name the input file (symlinks to it are also rejected).
`--target`	Explicit target. Defaults to `c`.

On success the CLI exits with 0 and writes the source to --output_file. On any failure it exits with 1, prints a single tilelang compile-only error: ... line to stderr, and leaves no stale output behind: a previous artifact at --output_file is removed before compilation, so a failed rerun cannot leave an earlier run’s source for a consumer such as Compiler Explorer to read.

Targets¶

Targets must be explicit. auto performs device detection and is therefore rejected, as a plain string and in JSON form:

$ python -I -m tilelang.tools.compile_only --target auto --output_file out.c example.py
tilelang compile-only error: target must be explicit; do not use auto

`--target` value	Behavior
`c` (default)	CPU C source. Works on every wheel, no GPU needed.
`cuda`	CUDA source pinned to `sm_80`, so no device detection happens.
`cuda -arch=sm_90`	CUDA source for the given architecture. Options must look like `-key=value`.
`{"kind": "cuda", "arch": "sm_90"}`	JSON target. A bare `{"kind": "cuda"}` is pinned to `sm_80`; `{"kind": "auto"}` is rejected.
`auto`	Rejected: `target must be explicit; do not use auto`.

On wheels without the CUDA codegen FFI (for example macOS Metal wheels), every CUDA target form fails softly with tilelang compile-only error: CUDA codegen FFI missing (e.g. macOS Metal wheel); use default --target c.

Error Reporting¶

Compile problems are reported as diagnostics on stderr with exit code 1, not as tracebacks about missing GPUs or drivers:

$ python -I -m tilelang.tools.compile_only --output_file out.c broken.py
tilelang compile-only error: no @tilelang.jit kernel or PrimFunc found

Syntax errors in the input, missing input files, and empty kernels are reported the same way.

Line Directives¶

When the installed build registers the tl.emit_line_directives pass config (PassConfigKey.TL_EMIT_LINE_DIRECTIVES), the tool enables it automatically and the generated source carries #line <n> "<input_file>" directives that map each statement back to the originating line of the example. Compiler Explorer uses these to link generated source to the input pane. Older wheels without the config key simply produce source without #line directives; no flag is needed either way.

Programmatic API¶

The same functionality is available in Python:

from tilelang.tools.compile_only import compile_kernel_source

source = compile_kernel_source(add.get_tir())              # default target "c"
cuda_source = compile_kernel_source(add.get_tir(), "cuda")  # pinned to sm_80

compile_kernel_source(func, target="c") lowers a PrimFunc (for example from JITImpl.get_tir()) and returns the non-empty kernel source string.
resolve_target(target) maps a CLI-style target string to the explicit target passed to lower(), applying the rules from the table above.
cuda_codegen_available() reports whether the installed wheel can lower CUDA targets.

Limitations¶

Only the first @tilelang.jit kernel or PrimFunc in the module is compiled; additional kernels in the same file are ignored.
The output is lowered kernel source, not a compiled binary: nothing is passed to nvcc, and PTX/SASS is out of scope.
The input file is executed with the invoking Python interpreter. Treat it like any script you run: python -I isolates it from your environment but does not sandbox it.