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)

tilelang_callback_hip_compile(code, target)

extrac_params(func)

canon_target_host(target, target_host)

host_codegen(host_mod, target_host)

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)¶
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 (Union[str, tvm.target.Target])

  • target_host (Optional[Union[str, tvm.target.Target]])

tilelang.engine.lower.host_codegen(host_mod, target_host)¶
Parameters:

  • host_mod (tilelang.tvm.IRModule)

  • target_host (tvm.target.Target)

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 (Union[tvm.tir.PrimFunc, tilelang.tvm.IRModule])

  • target (Union[str, tvm.target.Target])

  • target_host (Optional[Union[str, tvm.target.Target]])

Return type:

tilelang.engine.param.CompiledArtifact