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])

extrac_params(func)

host_codegen(host_mod, context)

Generate host-side code from the lowered IR module.

device_codegen(device_mod, context)

device_codegen_without_compile(device_mod, context)

lower_to_host_device_ir(func_or_mod, context[, ...])

Lower to split IR under a per-compilation instrumentation session.

lower_with_context(func_or_mod, context, *[, ...])

Compile using an existing context without resolving backend state again.

lower(func_or_mod[, target, target_host, ...])

Compile after resolving exactly one backend context.

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.tirx.PrimFunc)

tilelang.engine.lower.is_device_call(func)¶
Parameters:

func (tvm.tirx.PrimFunc)

tilelang.engine.lower.get_device_call(is_device_c=False)¶
Parameters:

is_device_c (bool)

Return type:

collections.abc.Callable[[tvm.tirx.PrimFunc], bool]

tilelang.engine.lower.get_host_call(is_device_c=False)¶
Parameters:

is_device_c (bool)

Return type:

collections.abc.Callable[[tvm.tirx.PrimFunc], bool]

tilelang.engine.lower.extrac_params(func)¶
Parameters:

func (tvm.tirx.PrimFunc)

Return type:

list[tilelang.engine.param.KernelParam]

tilelang.engine.lower.host_codegen(host_mod, context)¶

Generate host-side code from the lowered IR module.

Parameters:

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

  • context (BackendContext) – Resolved backend, device target, host target, and execution policy.

Return type:

tilelang.tvm.IRModule

tilelang.engine.lower.device_codegen(device_mod, context)¶
Parameters:
Return type:

tilelang.tvm.IRModule

tilelang.engine.lower.device_codegen_without_compile(device_mod, context)¶
Parameters:
Return type:

tilelang.tvm.IRModule

tilelang.engine.lower.lower_to_host_device_ir(func_or_mod, context, runtime_only=False)¶

Lower to split IR under a per-compilation instrumentation session.

Parameters:
Return type:

tuple[tilelang.tvm.IRModule, tilelang.tvm.IRModule, list[tilelang.engine.param.KernelParam] | None, tvm.target.Target, tvm.target.Target]

tilelang.engine.lower.lower_with_context(func_or_mod, context, *, runtime_only=False, enable_host_codegen=False, enable_device_compile=False)¶

Compile using an existing context without resolving backend state again.

enable_host_codegen builds the host runtime module when true. It is disabled by default because JIT adapters provide their own host-side integration. enable_device_compile similarly controls whether device code is compiled at this stage; JIT adapters compile it separately by default.

Parameters:
Return type:

tilelang.engine.param.CompiledArtifact

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

Compile after resolving exactly one backend context.

enable_host_codegen builds the host runtime module when true. It is disabled by default because JIT adapters provide their own host-side integration. enable_device_compile similarly controls whether device code is compiled at this stage; JIT adapters compile it separately by default.

Parameters:

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

  • target (str | tvm.target.Target)

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

  • runtime_only (bool)

  • enable_host_codegen (bool)

  • enable_device_compile (bool)

Return type:

tilelang.engine.param.CompiledArtifact