tilelang.jit.adapter.pto.adapter¶

Compile and execute PTODSL kernels using the PTO launch wrapper.

Classes¶

PTOKernelAdapter

Own PTO compilation and cached-library loading with a Cython launch ABI.

Functions¶

is_symbolic_expr(expr)

Check if the expression is a symbolic expression.

Module Contents¶

tilelang.jit.adapter.pto.adapter.is_symbolic_expr(expr)¶

Check if the expression is a symbolic expression. A symbolic expression can be a simple tvm.Var, or an tvm.PrimExpr containing tvm.Var.

Return type:

bool

class tilelang.jit.adapter.pto.adapter.PTOKernelAdapter(params, result_idx, target, func_or_mod, host_mod=None, device_mod=None, device_kernel_source=None, verbose=False, pass_configs=None, compile_flags=None)¶

Bases: tilelang.jit.adapter.base.BaseKernelAdapter

Own PTO compilation and cached-library loading with a Cython launch ABI.

Parameters:

  • params (list[tilelang.engine.param.KernelParam])

  • result_idx (list[int] | int | None)

  • target (str | dict[str, object] | tvm.target.Target)

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

  • host_mod (tilelang.tvm.IRModule | None)

  • device_mod (tilelang.tvm.IRModule | None)

  • device_kernel_source (str | None)

  • verbose (bool)

  • pass_configs (dict[str, Any] | None)

  • compile_flags (list[str] | None)

device_kernel_source = None¶
kernel_global_source = None¶
wrapper¶
host_kernel_source¶
classmethod from_database(params, result_idx, target, func_or_mod, host_kernel_source, device_kernel_source, kernel_lib_path, verbose=False, pass_configs=None, compile_flags=None)¶

Load the compiled PTO library without regenerating device or host code.

Parameters:
property prim_func: tvm.tirx.PrimFunc¶

Returns the primary TIR function from the IR module.

Return type:

tvm.tirx.PrimFunc

property srcpath¶

Returns the source path of the compiled library.

property libpath¶

Returns the path to the compiled library.

property lib_code¶

Returns the code of the compiled library.

property is_dynamic¶

Indicates whether the kernel handles dynamic shapes.

get_kernel_source(kernel_only=False)¶

Returns the source code of the compiled kernel.

Parameters:

kernel_only (bool)

get_host_source()¶

Returns the source code of the host function.