Skip to content

[tvm-ffi] Enable tvm-ffi for metal backend #1289

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Draft
wants to merge 2 commits into
base: main
Choose a base branch
from
Draft

[tvm-ffi] Enable tvm-ffi for metal backend #1289

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
ef2b518
upd
oraluben Nov 18, 2025
ab9ad61
enable tvm-ffi for metal
oraluben Nov 19, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
10 changes: 2 additions & 8 deletions CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -260,15 +260,9 @@ install(TARGETS tilelang_cython_wrapper

# let libtilelang to search tvm/tvm_runtime in same dir
if(APPLE)
set_target_properties(tilelang PROPERTIES INSTALL_RPATH "@loader_path;@loader_path/../../tvm_ffi/lib")
set_target_properties(tilelang_module PROPERTIES INSTALL_RPATH "@loader_path;@loader_path/../../tvm_ffi/lib")
set_target_properties(tvm PROPERTIES INSTALL_RPATH "@loader_path;@loader_path/../../tvm_ffi/lib")
set_target_properties(tvm_runtime PROPERTIES INSTALL_RPATH "@loader_path;@loader_path/../../tvm_ffi/lib")
set_property(GLOBAL PROPERTY INSTALL_RPATH "@loader_path;@loader_path/../../tvm_ffi/lib")
elseif(UNIX)
set_target_properties(tilelang PROPERTIES INSTALL_RPATH "\$ORIGIN:\$ORIGIN/../../tvm_ffi/lib")
set_target_properties(tilelang_module PROPERTIES INSTALL_RPATH "\$ORIGIN:\$ORIGIN/../../tvm_ffi/lib")
set_target_properties(tvm PROPERTIES INSTALL_RPATH "\$ORIGIN:\$ORIGIN/../../tvm_ffi/lib")
set_target_properties(tvm_runtime PROPERTIES INSTALL_RPATH "\$ORIGIN:\$ORIGIN/../../tvm_ffi/lib")
set_property(GLOBAL PROPERTY INSTALL_RPATH "\$ORIGIN:\$ORIGIN/../../tvm_ffi/lib")
endif()

install(
Expand Down
2 changes: 1 addition & 1 deletion pyproject.toml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -27,7 +27,7 @@ classifiers = [
]
dynamic = ["version"]
dependencies = [
"apache-tvm-ffi==0.1.0",
"apache-tvm-ffi~=0.1.0",
"cloudpickle",
"ml-dtypes",
"numpy>=1.23.5",
Expand Down
2 changes: 1 addition & 1 deletion testing/python/metal/test_metal_codegen.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,7 +5,7 @@
import torch


@tilelang.jit(execution_backend='torch')
@tilelang.jit(execution_backend='tvm_ffi')
def matmul(M, N, K, block_M, block_N, block_K, dtype="float32", accum_dtype="float"):

@T.prim_func
Expand Down
2 changes: 2 additions & 0 deletions tilelang/engine/lower.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -180,6 +180,8 @@ def device_codegen(device_mod: tvm.IRModule, target: Target) -> tvm.IRModule:
device_mod = tvm.ffi.get_global_func("target.build.tilelang_cuda")(device_mod, target)
elif target.kind.name == "hip":
device_mod = tvm.ffi.get_global_func("target.build.tilelang_hip")(device_mod, target)
elif target.kind.name == "metal":
device_mod = tvm.ffi.get_global_func("target.build.metal")(device_mod, target)
else:
raise ValueError(f"Target {target.kind.name} is not supported")

Expand Down
2 changes: 1 addition & 1 deletion tilelang/jit/execution_backend.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -35,7 +35,7 @@ def allowed_backends_for_target(target: Target, *, include_unavailable: bool = T
elif kind == "hip":
allowed = ["tvm_ffi", "cython", "ctypes"]
elif kind == "metal":
allowed = ["torch"]
allowed = ["tvm_ffi", "torch"]
elif kind == "c": # CPU C backend
allowed = ["cython", "ctypes", "tvm_ffi"]
else:
Expand Down