tilelang.engine.lower

The compiler for TL programs.

Functions

is_cpu_device_backend(target)
has_device_kernel_launch(attrs)	Check if the attributes indicate a device kernel launch.
is_device_call_c_device(func)
is_device_call(func)
get_device_call([is_device_c])
get_host_call([is_device_c])
tilelang_callback_cuda_compile(code, target[, pass_config])
tilelang_callback_hip_compile(code, target)
extrac_params(func)
canon_target_host(target, target_host)
host_codegen(host_mod, target_host[, target])	Generate host-side code from the lowered IR module.
device_codegen(device_mod, target)
device_codegen_without_compile(device_mod, target)
lower(func_or_mod[, target, target_host, ...])	enable_host_codegen: whether to enable host codegen, default is False, as we have our

Module Contents

tilelang.engine.lower.is_cpu_device_backend(target)

Parameters:

target (tvm.target.Target)

tilelang.engine.lower.has_device_kernel_launch(attrs)

Check if the attributes indicate a device kernel launch.

Return type:

bool

tilelang.engine.lower.is_device_call_c_device(func)

Parameters:

func (tvm.tir.PrimFunc)

tilelang.engine.lower.is_device_call(func)

Parameters:

func (tvm.tir.PrimFunc)

tilelang.engine.lower.get_device_call(is_device_c=False)

Parameters:

is_device_c (bool)

Return type:

Callable[[tvm.tir.PrimFunc], bool]

tilelang.engine.lower.get_host_call(is_device_c=False)

Parameters:

is_device_c (bool)

Return type:

Callable[[tvm.tir.PrimFunc], bool]

tilelang.engine.lower.tilelang_callback_cuda_compile(code, target, pass_config=None)

tilelang.engine.lower.tilelang_callback_hip_compile(code, target)

tilelang.engine.lower.extrac_params(func)

Parameters:

func (tvm.tir.PrimFunc)

Return type:

list[tilelang.engine.param.KernelParam]

tilelang.engine.lower.canon_target_host(target, target_host)

Parameters:

• target (str | tvm.target.Target)

• target_host (str | tvm.target.Target | None)

tilelang.engine.lower.host_codegen(host_mod, target_host, target=None)

Generate host-side code from the lowered IR module.

Parameters:

• host_mod (tvm.IRModule) – The host-side IR module to compile.

• target_host (Target) – The host compilation target (e.g. “llvm” or “c”).

• target (Target, optional) – The device target. When the device target is Metal, the pass MarkHostMetalContext is applied so that the generated host code contains the Metal/MPS synchronisation logic.

Return type:

tilelang.tvm.IRModule

tilelang.engine.lower.device_codegen(device_mod, target)

Parameters:

• device_mod (tilelang.tvm.IRModule)

• target (tvm.target.Target)

Return type:

tilelang.tvm.IRModule

tilelang.engine.lower.device_codegen_without_compile(device_mod, target)

Parameters:

• device_mod (tilelang.tvm.IRModule)

• target (tvm.target.Target)

Return type:

tilelang.tvm.IRModule

tilelang.engine.lower.lower(func_or_mod, target='auto', target_host=None, runtime_only=False, enable_host_codegen=False, enable_device_compile=False)

enable_host_codegen: whether to enable host codegen, default is False, as we have our own host codegen implementation in jit. enable_device_compile: whether to enable device codegen, default is False, as we have our own device codegen implementation in jit.

Parameters:

• func_or_mod (tvm.tir.PrimFunc | tilelang.tvm.IRModule)

• target (str | tvm.target.Target)

• target_host (str | tvm.target.Target | None)

Return type:

tilelang.engine.param.CompiledArtifact