You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by tq...@apache.org on 2025/05/05 14:36:54 UTC
(tvm) branch refactor-s3 updated (251e89e811 -> b5d5d6c398)
This is an automated email from the ASF dual-hosted git repository.
tqchen pushed a change to branch refactor-s3
in repository https://gitbox.apache.org/repos/asf/tvm.git
discard 251e89e811 [WIN] Remove lock to avoid mutex dep at init time
discard 5438ef692e Fix C codegen caller convention
discard 6e85d026e6 [WIN] Cleanup for win
discard 84d40373bf [FFI] Simplify unpack traits
discard 3aaf85cbd7 Fix through relax
new 528b0d1c31 Fix through relax
new 0ddf789068 [FFI] Simplify unpack traits
new a81a3f313b [CYTHON] Phase out legacy FFI
new 0c7211ec1e [WIN] Cleanup for win
new e949a9d9c1 Fix C codegen caller convention
new b5d5d6c398 [WIN] Remove lock to avoid mutex dep at init time
This update added new revisions after undoing existing revisions.
That is to say, some revisions that were in the old version of the
branch are not in the new version. This situation occurs
when a user --force pushes a change and generates a repository
containing something like this:
* -- * -- B -- O -- O -- O (251e89e811)
\
N -- N -- N refs/heads/refactor-s3 (b5d5d6c398)
You should already have received notification emails for all of the O
revisions, and so the following emails describe only the N revisions
from the common base, B.
Any revisions marked "omit" are not gone; other references still
refer to them. Any revisions marked "discard" are gone forever.
The 6 revisions listed above as "new" are entirely new to this
repository and will be described in separate emails. The revisions
listed as "add" were already present in the repository and have only
been added to this reference.
Summary of changes:
ffi/include/tvm/ffi/function_details.h | 3 +-
python/setup.py | 21 --
python/tvm/_ffi/_cy3/__init__.py | 17 --
python/tvm/_ffi/_cython/base.pxi | 236 --------------------
python/tvm/_ffi/_cython/core.pyx | 21 --
python/tvm/_ffi/_cython/ndarray.pxi | 180 ---------------
python/tvm/_ffi/_cython/object.pxi | 165 --------------
python/tvm/_ffi/_cython/packed_func.pxi | 376 --------------------------------
python/tvm/runtime/ndarray.py | 2 +-
src/runtime/metal/metal_module.mm | 2 +-
src/tir/ir/expr.cc | 55 ++---
11 files changed, 32 insertions(+), 1046 deletions(-)
delete mode 100644 python/tvm/_ffi/_cy3/__init__.py
delete mode 100644 python/tvm/_ffi/_cython/base.pxi
delete mode 100644 python/tvm/_ffi/_cython/core.pyx
delete mode 100644 python/tvm/_ffi/_cython/ndarray.pxi
delete mode 100644 python/tvm/_ffi/_cython/object.pxi
delete mode 100644 python/tvm/_ffi/_cython/packed_func.pxi
(tvm) 01/06: Fix through relax
Posted by tq...@apache.org.
This is an automated email from the ASF dual-hosted git repository.
tqchen pushed a commit to branch refactor-s3
in repository https://gitbox.apache.org/repos/asf/tvm.git
commit 528b0d1c3109d347b0568b27a85b0120acc0ab8b
Author: tqchen <ti...@gmail.com>
AuthorDate: Sun May 4 19:41:53 2025 -0400
Fix through relax
---
python/tvm/ffi/container.py | 3 ++
python/tvm/ffi/cython/dtype.pxi | 2 +-
python/tvm/ffi/cython/object.pxi | 3 +-
python/tvm/ffi/dtype.py | 21 +++++++++
python/tvm/meta_schedule/cost_model/mlp_model.py | 5 +-
python/tvm/relax/op/create.py | 2 +-
python/tvm/runtime/ndarray.py | 6 +--
.../test_runtime_packed_func.py | 1 -
.../contrib/test_hexagon/test_parallel_scalar.py | 6 +--
tests/python/ffi/test_dtype.py | 8 ++++
tests/python/ffi/test_string.py | 17 +++++++
tests/python/relax/test_op_inspect.py | 53 +++++-----------------
tests/python/relax/test_op_misc.py | 2 +-
tests/python/relax/test_vm_build.py | 8 ++--
tests/python/relax/test_vm_callback_function.py | 11 ++---
tests/python/relax/test_vm_codegen_only.py | 2 +-
16 files changed, 80 insertions(+), 70 deletions(-)
diff --git a/python/tvm/ffi/container.py b/python/tvm/ffi/container.py
index ddddd2d7cc..6ababe2557 100644
--- a/python/tvm/ffi/container.py
+++ b/python/tvm/ffi/container.py
@@ -100,6 +100,9 @@ class KeysView(collections.abc.KeysView):
if not functor(2):
break
+ def __contains__(self, k):
+ return self._backend_map.__contains__(k)
+
class ValuesView(collections.abc.ValuesView):
"""Helper class to return values view"""
diff --git a/python/tvm/ffi/cython/dtype.pxi b/python/tvm/ffi/cython/dtype.pxi
index bbf9e60053..ef71ea4edd 100644
--- a/python/tvm/ffi/cython/dtype.pxi
+++ b/python/tvm/ffi/cython/dtype.pxi
@@ -28,7 +28,7 @@ def _create_dtype_from_tuple(cls, code, bits, lanes):
cdtype.code = code
cdtype.bits = bits
cdtype.lanes = lanes
- ret = cls.__new__(cls)
+ ret = cls.__new__(cls, str(cdtype))
(<DataType>ret).cdtype = cdtype
return ret
diff --git a/python/tvm/ffi/cython/object.pxi b/python/tvm/ffi/cython/object.pxi
index 1ac32c3bc6..c258f578b0 100644
--- a/python/tvm/ffi/cython/object.pxi
+++ b/python/tvm/ffi/cython/object.pxi
@@ -125,7 +125,8 @@ cdef class Object:
return __object_dir__(self)
def __repr__(self):
- return __object_repr__(self)
+ # make sure repr is a raw string
+ return str(__object_repr__(self))
def __eq__(self, other):
return self.same_as(other)
diff --git a/python/tvm/ffi/dtype.py b/python/tvm/ffi/dtype.py
index 56b888316d..32986a4eb0 100644
--- a/python/tvm/ffi/dtype.py
+++ b/python/tvm/ffi/dtype.py
@@ -84,6 +84,27 @@ class dtype(str):
def __repr__(self):
return f"dtype('{self}')"
+ def with_lanes(self, lanes):
+ """
+ Create a new dtype with the given number of lanes.
+
+ Parameters
+ ----------
+ lanes : int
+ The number of lanes.
+
+ Returns
+ -------
+ dtype
+ The new dtype with the given number of lanes.
+ """
+ cdtype = core._create_dtype_from_tuple(
+ core.DataType, self.__tvm_ffi_dtype__.type_code, self.__tvm_ffi_dtype__.bits, lanes
+ )
+ val = str.__new__(dtype, str(cdtype))
+ val.__tvm_ffi_dtype__ = cdtype
+ return val
+
@property
def itemsize(self):
return self.__tvm_ffi_dtype__.itemsize
diff --git a/python/tvm/meta_schedule/cost_model/mlp_model.py b/python/tvm/meta_schedule/cost_model/mlp_model.py
index 4ee5ba838d..9167d30e90 100644
--- a/python/tvm/meta_schedule/cost_model/mlp_model.py
+++ b/python/tvm/meta_schedule/cost_model/mlp_model.py
@@ -534,6 +534,7 @@ class State:
if json_file.endswith("_workload.json"):
workload_paths.append(json_file)
for workload_path in tqdm(workload_paths):
+ # pylint: disable=protected-access,broad-exception-caught
try:
database = JSONDatabase(
path_workload=workload_path,
@@ -541,9 +542,7 @@ class State:
"_workload.json", "_candidates.json"
),
)
- except (
- tvm._ffi.base.TVMError
- ): # pylint: disable=protected-access,broad-exception-caught
+ except tvm._ffi.base.TVMError:
continue
candidates, results = [], []
tuning_records = database.get_all_tuning_records()
diff --git a/python/tvm/relax/op/create.py b/python/tvm/relax/op/create.py
index 029682ded8..c61d9521a4 100644
--- a/python/tvm/relax/op/create.py
+++ b/python/tvm/relax/op/create.py
@@ -17,7 +17,7 @@
"""Creation operators."""
from typing import Optional, Tuple, Union
-from tvm import DataType
+from tvm import DataType, DataTypeCode
from tvm.ir.expr import PrimExpr
from . import _ffi_api
diff --git a/python/tvm/runtime/ndarray.py b/python/tvm/runtime/ndarray.py
index 5581adbbc1..7c29805092 100644
--- a/python/tvm/runtime/ndarray.py
+++ b/python/tvm/runtime/ndarray.py
@@ -92,7 +92,7 @@ class NDArray(tvm.ffi.core.NDArray):
):
raise ValueError("Array only support set from numpy array")
if isinstance(value, NDArray):
- if value.handle is not self.handle:
+ if not value.same_as(self):
value.copyto(self)
elif isinstance(value, (np.ndarray, np.generic)):
self.copyfrom(value)
@@ -128,7 +128,7 @@ class NDArray(tvm.ffi.core.NDArray):
shape, dtype = self.shape, self.dtype
if t.lanes > 1:
shape = shape + (t.lanes,)
- t.lanes = 1
+ t = t.with_lanes(1)
dtype = str(t)
if source_array.shape != shape:
@@ -185,7 +185,7 @@ class NDArray(tvm.ffi.core.NDArray):
old_dtype = dtype
if t.lanes > 1:
shape = shape + (t.lanes,)
- t.lanes = 1
+ t = t.with_lanes(1)
dtype = str(t)
if dtype == "int4":
dtype = "int8"
diff --git a/tests/python/all-platform-minimal-test/test_runtime_packed_func.py b/tests/python/all-platform-minimal-test/test_runtime_packed_func.py
index f24e1edde7..f315b8f3c2 100644
--- a/tests/python/all-platform-minimal-test/test_runtime_packed_func.py
+++ b/tests/python/all-platform-minimal-test/test_runtime_packed_func.py
@@ -45,7 +45,6 @@ def test_get_callback_with_node():
x = T.int32(10)
def test(y):
- assert y.handle != x.handle
return y
f2 = tvm.runtime.convert(test)
diff --git a/tests/python/contrib/test_hexagon/test_parallel_scalar.py b/tests/python/contrib/test_hexagon/test_parallel_scalar.py
index d7fd5a3b20..bd9c78d5da 100644
--- a/tests/python/contrib/test_hexagon/test_parallel_scalar.py
+++ b/tests/python/contrib/test_hexagon/test_parallel_scalar.py
@@ -117,11 +117,7 @@ def evaluate(hexagon_session, operations, expected, sch):
class TestMatMulVec:
"""MatMul test class."""
- (
- operation_name,
- operator_producer,
- expected_output_producer,
- ) = tvm.testing.parameters(
+ (operation_name, operator_producer, expected_output_producer,) = tvm.testing.parameters(
("add", get_add_operator, (lambda a, b: a + b)),
("mul", get_multiply_operator, (lambda a, b: a * b)),
("sub", get_sub_operator, (lambda a, b: a - b)),
diff --git a/tests/python/ffi/test_dtype.py b/tests/python/ffi/test_dtype.py
index fb6e14a17a..2758edf9d6 100644
--- a/tests/python/ffi/test_dtype.py
+++ b/tests/python/ffi/test_dtype.py
@@ -54,3 +54,11 @@ def test_dtype_pickle(dtype_str):
assert dtype_pickled.type_code == dtype.type_code
assert dtype_pickled.bits == dtype.bits
assert dtype_pickled.lanes == dtype.lanes
+
+
+def test_dtype_with_lanes():
+ dtype = tvm_ffi.dtype("float32")
+ dtype_with_lanes = dtype.with_lanes(4)
+ assert dtype_with_lanes.type_code == dtype.type_code
+ assert dtype_with_lanes.bits == dtype.bits
+ assert dtype_with_lanes.lanes == 4
diff --git a/tests/python/ffi/test_string.py b/tests/python/ffi/test_string.py
index 98eab5bcb7..cac948b53d 100644
--- a/tests/python/ffi/test_string.py
+++ b/tests/python/ffi/test_string.py
@@ -1,3 +1,20 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+
import pickle
from tvm import ffi as tvm_ffi
diff --git a/tests/python/relax/test_op_inspect.py b/tests/python/relax/test_op_inspect.py
index f8326db8dd..2ba9f9a709 100644
--- a/tests/python/relax/test_op_inspect.py
+++ b/tests/python/relax/test_op_inspect.py
@@ -158,10 +158,6 @@ def test_strides_of_compact_tensor(shape):
def test_strides_of_non_compact_tensor():
- backing_shape = [64, 64]
- view_shape = [16, 16]
- expected_strides = [backing_shape[0], 1]
-
@I.ir_module
class mod:
@R.function
@@ -170,20 +166,12 @@ def test_strides_of_non_compact_tensor():
built = tvm.compile(mod)
vm = relax.VirtualMachine(built, tvm.cpu())
-
- backing_ndarray = tvm.nd.empty(backing_shape, "int32")
-
- # Manually overwrite the DLTensor fields to make a view into the
- # tensor.
- view = backing_ndarray.handle[0]
- np_shape = np.array([16, 16], "int64")
- view.shape = np_shape.ctypes.data_as(ctypes.POINTER(ctypes.c_long))
- np_strides = np.array([64, 1], "int64")
- view.strides = np_strides.ctypes.data_as(ctypes.POINTER(ctypes.c_long))
- backing_ndarray.handle[0] = view
-
- res = [vm["main"](backing_ndarray, i) for i, _ in enumerate(view_shape)]
-
+ view_shape = [4, 4]
+ expected_strides = [1, 4]
+ # use transpose to make strides non-compact
+ x = np.zeros([4, 4], "int32").T
+ y = tvm.ffi.from_dlpack(x, required_alignment=4, required_contiguous=False)
+ res = [vm["main"](y, i) for i, _ in enumerate(view_shape)]
tvm.ir.assert_structural_equal(res, expected_strides)
@@ -200,19 +188,10 @@ def test_byte_offset(elem_offset):
built = tvm.compile(mod)
vm = relax.VirtualMachine(built, tvm.cpu())
-
- backing_ndarray = tvm.nd.empty(backing_shape, "int32")
-
- # Manually overwrite the DLTensor fields to make a view into the
- # tensor.
- view = backing_ndarray.handle[0]
- np_shape = np.array(view_shape, "int64")
- view.shape = np_shape.ctypes.data_as(ctypes.POINTER(ctypes.c_long))
- view.byte_offset = byte_offset
- backing_ndarray.handle[0] = view
-
- res = vm["main"](backing_ndarray)
-
+ dtype = "int32"
+ backing_ndarray = tvm.nd.empty(backing_shape, dtype)
+ view = backing_ndarray._create_view(view_shape, dtype, relative_byte_offset=byte_offset)
+ res = vm["main"](view)
assert res == byte_offset
@@ -234,16 +213,8 @@ def test_elem_offset(elem_offset, dtype):
vm = relax.VirtualMachine(built, tvm.cpu())
backing_ndarray = tvm.nd.empty(backing_shape, dtype)
-
- # Manually overwrite the DLTensor fields to make a view into the
- # tensor.
- view = backing_ndarray.handle[0]
- np_shape = np.array(view_shape, "int64")
- view.shape = np_shape.ctypes.data_as(ctypes.POINTER(ctypes.c_long))
- view.byte_offset = byte_offset
- backing_ndarray.handle[0] = view
-
- res = vm["main"](backing_ndarray)
+ view = backing_ndarray._create_view(view_shape, dtype, relative_byte_offset=byte_offset)
+ res = vm["main"](view)
assert res == elem_offset
diff --git a/tests/python/relax/test_op_misc.py b/tests/python/relax/test_op_misc.py
index b53c10b369..366ea1b688 100644
--- a/tests/python/relax/test_op_misc.py
+++ b/tests/python/relax/test_op_misc.py
@@ -58,7 +58,7 @@ def test_call_tir_with_grad():
)
assert v2.attrs.te_grad_name == "identity_k_grad"
assert isinstance(v2.attrs.te_grad_kwargs, tvm.ir.container.Map)
- val = v2.attrs.te_grad_kwargs.items()[0]
+ val = list(v2.attrs.te_grad_kwargs.items())[0]
assert val[0] == "k" and float(val[1]) == 1.0
diff --git a/tests/python/relax/test_vm_build.py b/tests/python/relax/test_vm_build.py
index 9acd1b8629..2078248880 100644
--- a/tests/python/relax/test_vm_build.py
+++ b/tests/python/relax/test_vm_build.py
@@ -541,7 +541,7 @@ def test_vm_relax_symbolic_shape_tuple(exec_mode):
func = vm["main"]
- assert func(ShapeTuple([2, 3])) == [4, 9]
+ assert func(ShapeTuple([2, 3])) == (4, 9)
with pytest.raises(ValueError):
func(ShapeTuple([2, 3, 4]))
@@ -595,7 +595,7 @@ def test_vm_relax_multiple_symbolic_prim_value(exec_mode):
func = vm["main"]
- assert func(2, ShapeTuple([4, 12]), 6) == [4, 7]
+ assert func(2, ShapeTuple([4, 12]), 6) == (4, 7)
with pytest.raises(RuntimeError):
func(2, ShapeTuple([4, 12]), 1)
@@ -873,8 +873,8 @@ def test_vm_to_device(exec_mode):
res_2 = check_saved_func(vm, "foo2", x_inp)
# check the copied tensor's device
- assert str(res_1.device) == "cuda(0)"
- assert str(res_2.device) == "cpu(0)"
+ assert res_1.device == tvm.cuda(0)
+ assert res_2.device == tvm.cpu(0)
tvm.testing.assert_allclose(res_1.numpy(), x_inp.numpy())
tvm.testing.assert_allclose(res_2.numpy(), x_inp.numpy())
diff --git a/tests/python/relax/test_vm_callback_function.py b/tests/python/relax/test_vm_callback_function.py
index 73336db559..1cee0b57d8 100644
--- a/tests/python/relax/test_vm_callback_function.py
+++ b/tests/python/relax/test_vm_callback_function.py
@@ -111,15 +111,10 @@ def test_catch_exception_with_full_stack_trace(exec_mode, target, dev):
while stack.tb_next is not None:
stack = stack.tb_next
frame = stack.tb_frame
+ assert (
+ frame.f_code.co_filename.find("test_vm_callback_function.py") != -1
+ ), "Inner-most stack frame should be from Python callback"
- assert frame.f_code is custom_callback.__code__, (
- "Inner-most stack frame should be from Python callback, "
- "even though that crosses an FFI boundary"
- )
- assert frame.f_locals.get("local_var") == 42, (
- "Python __traceback__ should include local variables, "
- "even though that crosses an FFI boundary"
- )
else:
raise RuntimeError("Exception thrown in callback was not propagated to calling scope")
diff --git a/tests/python/relax/test_vm_codegen_only.py b/tests/python/relax/test_vm_codegen_only.py
index 5ae46b099e..dac0f867ce 100644
--- a/tests/python/relax/test_vm_codegen_only.py
+++ b/tests/python/relax/test_vm_codegen_only.py
@@ -78,7 +78,7 @@ def test_vm_to_device(exec_mode):
res = check_saved_func(vm, "foo", inp)
tvm.testing.assert_allclose(res.numpy(), inp.numpy(), rtol=1e-7, atol=1e-7)
# check the resulting tensor is on cpu:0
- assert str(res.device) == "cpu(0)"
+ assert res.device == tvm.cpu(0)
assert res.device.device_type == 1
assert res.device.device_id == 0
(tvm) 06/06: [WIN] Remove lock to avoid mutex dep at init time
Posted by tq...@apache.org.
This is an automated email from the ASF dual-hosted git repository.
tqchen pushed a commit to branch refactor-s3
in repository https://gitbox.apache.org/repos/asf/tvm.git
commit b5d5d6c398cbc67dec03d25debbb3be39713a574
Author: tqchen <ti...@gmail.com>
AuthorDate: Mon May 5 09:35:15 2025 -0400
[WIN] Remove lock to avoid mutex dep at init time
---
src/node/attr_registry.h | 7 -------
1 file changed, 7 deletions(-)
diff --git a/src/node/attr_registry.h b/src/node/attr_registry.h
index 2b022d3e45..12913ccefd 100644
--- a/src/node/attr_registry.h
+++ b/src/node/attr_registry.h
@@ -28,7 +28,6 @@
#include <tvm/runtime/packed_func.h>
#include <memory>
-#include <mutex>
#include <unordered_map>
#include <utility>
#include <vector>
@@ -95,7 +94,6 @@ class AttrRegistry {
*/
void UpdateAttr(const String& attr_name, const KeyType& key, Any value, int plevel) {
using runtime::TVMRetValue;
- std::lock_guard<std::mutex> lock(mutex_);
auto& op_map = attrs_[attr_name];
if (op_map == nullptr) {
op_map.reset(new AttrRegistryMapContainerMap<KeyType>());
@@ -122,7 +120,6 @@ class AttrRegistry {
* \param key The key to the attribute table.
*/
void ResetAttr(const String& attr_name, const KeyType& key) {
- std::lock_guard<std::mutex> lock(mutex_);
auto& op_map = attrs_[attr_name];
if (op_map == nullptr) {
return;
@@ -139,7 +136,6 @@ class AttrRegistry {
* \return The result attribute map.
*/
const AttrRegistryMapContainerMap<KeyType>& GetAttrMap(const String& attr_name) {
- std::lock_guard<std::mutex> lock(mutex_);
auto it = attrs_.find(attr_name);
if (it == attrs_.end()) {
LOG(FATAL) << "Attribute \'" << attr_name << "\' is not registered";
@@ -153,7 +149,6 @@ class AttrRegistry {
* \return The check result.
*/
bool HasAttrMap(const String& attr_name) {
- std::lock_guard<std::mutex> lock(mutex_);
return attrs_.count(attr_name);
}
@@ -166,8 +161,6 @@ class AttrRegistry {
}
private:
- // mutex to avoid registration from multiple threads.
- std::mutex mutex_;
// entries in the registry
std::vector<std::unique_ptr<EntryType>> entries_;
// map from name to entries.
(tvm) 04/06: [WIN] Cleanup for win
Posted by tq...@apache.org.
This is an automated email from the ASF dual-hosted git repository.
tqchen pushed a commit to branch refactor-s3
in repository https://gitbox.apache.org/repos/asf/tvm.git
commit 0c7211ec1e75bae7d4a79b286f72bcde778ae402
Author: tqchen <tq...@gmail.com>
AuthorDate: Sun May 4 21:20:06 2025 -0400
[WIN] Cleanup for win
---
cmake/utils/FindLLVM.cmake | 5 ++++-
ffi/CMakeLists.txt | 7 +++++++
ffi/src/ffi/traceback_win.cc | 3 +--
3 files changed, 12 insertions(+), 3 deletions(-)
diff --git a/cmake/utils/FindLLVM.cmake b/cmake/utils/FindLLVM.cmake
index 182a2c6693..2a243b06c8 100644
--- a/cmake/utils/FindLLVM.cmake
+++ b/cmake/utils/FindLLVM.cmake
@@ -199,7 +199,7 @@ macro(find_llvm use_llvm)
message(STATUS "LLVM links against zlib")
find_package(ZLIB REQUIRED)
list(APPEND LLVM_LIBS "ZLIB::ZLIB")
- elseif("${__flag}" STREQUAL "-lzstd" OR ("${__flag}" STREQUAL "zstd.dll.lib"))
+ elseif("${__flag}" STREQUAL "-lzstd")
list(APPEND CMAKE_MODULE_PATH "${__llvm_cmakedir}")
find_package(zstd REQUIRED)
if (TARGET "zstd::libzstd_static")
@@ -212,6 +212,9 @@ macro(find_llvm use_llvm)
elseif("${__flag}" STREQUAL "-lxml2")
message(STATUS "LLVM links against xml2")
list(APPEND LLVM_LIBS "-lxml2")
+ elseif("${__flag}" STREQUAL "zstd.dll.lib")
+ message(STATUS "LLVM linker flag under LLVM libdir: ${__llvm_libdir}/zstd.lib")
+ list(APPEND LLVM_LIBS "${__llvm_libdir}/zstd.lib")
elseif((__flag MATCHES ".lib$") AND (EXISTS "${__llvm_libdir}/${__flag}"))
# If the library file ends in .lib try to also search the llvm_libdir
message(STATUS "LLVM linker flag under LLVM libdir: ${__llvm_libdir}/${__flag}")
diff --git a/ffi/CMakeLists.txt b/ffi/CMakeLists.txt
index 209913d60f..ab9a3fabc3 100644
--- a/ffi/CMakeLists.txt
+++ b/ffi/CMakeLists.txt
@@ -104,6 +104,13 @@ if (TARGET libbacktrace)
target_link_libraries(tvm_ffi_shared PRIVATE libbacktrace)
target_link_libraries(tvm_ffi_static PRIVATE libbacktrace)
endif ()
+
+if (MSVC)
+ target_link_libraries(tvm_ffi_objs PRIVATE DbgHelp.lib)
+ target_link_libraries(tvm_ffi_shared PRIVATE DbgHelp.lib)
+ target_link_libraries(tvm_ffi_static PRIVATE DbgHelp.lib)
+endif ()
+
target_link_libraries(tvm_ffi_objs PUBLIC tvm_ffi_header)
target_link_libraries(tvm_ffi_shared PUBLIC tvm_ffi_header)
target_link_libraries(tvm_ffi_static PUBLIC tvm_ffi_header)
diff --git a/ffi/src/ffi/traceback_win.cc b/ffi/src/ffi/traceback_win.cc
index 166b7eb242..4c89c4f239 100644
--- a/ffi/src/ffi/traceback_win.cc
+++ b/ffi/src/ffi/traceback_win.cc
@@ -23,9 +23,8 @@
*/
#ifdef _MSC_VER
-#include <dbghelp.h>
#include <windows.h>
-#pragma comment(lib, "DbgHelp.lib")
+#include <dbghelp.h> // NOLINT(*)
#include <tvm/ffi/c_api.h>
#include <tvm/ffi/error.h>
(tvm) 05/06: Fix C codegen caller convention
Posted by tq...@apache.org.
This is an automated email from the ASF dual-hosted git repository.
tqchen pushed a commit to branch refactor-s3
in repository https://gitbox.apache.org/repos/asf/tvm.git
commit e949a9d9c12aaf0d9d4ad7e9af265b5d9987eb31
Author: tqchen <ti...@gmail.com>
AuthorDate: Mon May 5 09:24:27 2025 -0400
Fix C codegen caller convention
---
src/target/source/codegen_c_host.cc | 6 +++++-
1 file changed, 5 insertions(+), 1 deletion(-)
diff --git a/src/target/source/codegen_c_host.cc b/src/target/source/codegen_c_host.cc
index 815904c946..ef86f09ca2 100644
--- a/src/target/source/codegen_c_host.cc
+++ b/src/target/source/codegen_c_host.cc
@@ -242,7 +242,11 @@ void CodeGenCHost::PrintCallPacked(const CallNode* op) {
std::string result = name_supply_->FreshName("result");
this->stream << "TVMFFIAny " << result << ";\n";
this->PrintIndent();
-
+ // must make sure type_index is set to none
+ this->stream << result << ".type_index = kTVMFFINone;\n";
+ this->PrintIndent();
+ this->stream << result << ".v_int64 = 0;\n";
+ this->PrintIndent();
if (op->op.same_as(builtin::tvm_call_packed_lowered())) {
this->stream << "if (TVMFFIFunctionCall(" << packed_func_name << ", ";
} else {
(tvm) 03/06: [CYTHON] Phase out legacy FFI
Posted by tq...@apache.org.
This is an automated email from the ASF dual-hosted git repository.
tqchen pushed a commit to branch refactor-s3
in repository https://gitbox.apache.org/repos/asf/tvm.git
commit a81a3f313b3b7a894ae61d27f84ad51e90322502
Author: tqchen <tq...@gmail.com>
AuthorDate: Sun May 4 21:35:37 2025 -0400
[CYTHON] Phase out legacy FFI
---
python/setup.py | 21 --
python/tvm/_ffi/_cy3/__init__.py | 17 --
python/tvm/_ffi/_cython/base.pxi | 236 --------------------
python/tvm/_ffi/_cython/core.pyx | 21 --
python/tvm/_ffi/_cython/ndarray.pxi | 180 ---------------
python/tvm/_ffi/_cython/object.pxi | 165 --------------
python/tvm/_ffi/_cython/packed_func.pxi | 376 --------------------------------
7 files changed, 1016 deletions(-)
diff --git a/python/setup.py b/python/setup.py
index 02cac5f897..3900e5b3d0 100644
--- a/python/setup.py
+++ b/python/setup.py
@@ -148,7 +148,6 @@ def config_cython():
subdir = "_cy3"
ret = []
- cython_source = "tvm/_ffi/_cython"
extra_compile_args = ["-std=c++17", "-DDMLC_USE_LOGGING_LIBRARY=<tvm/runtime/logging.h>"]
if os.name == "nt":
library_dirs = ["tvm", "../build/Release", "../build"]
@@ -164,26 +163,6 @@ def config_cython():
library_dirs = None
libraries = None
- for fn in os.listdir(cython_source):
- if not fn.endswith(".pyx"):
- continue
- ret.append(
- Extension(
- "tvm._ffi.%s.%s" % (subdir, fn[:-4]),
- ["tvm/_ffi/_cython/%s" % fn],
- include_dirs=[
- "../include/",
- "../3rdparty/dmlc-core/include",
- "../ffi/include/",
- "../ffi/3rdparty/dlpack/include",
- ],
- extra_compile_args=extra_compile_args,
- library_dirs=library_dirs,
- libraries=libraries,
- language="c++",
- )
- )
-
# the latest ffi source
for fn in os.listdir("tvm/ffi/cython"):
if not fn.endswith(".pyx"):
diff --git a/python/tvm/_ffi/_cy3/__init__.py b/python/tvm/_ffi/_cy3/__init__.py
deleted file mode 100644
index 159f3254db..0000000000
--- a/python/tvm/_ffi/_cy3/__init__.py
+++ /dev/null
@@ -1,17 +0,0 @@
-# Licensed to the Apache Software Foundation (ASF) under one
-# or more contributor license agreements. See the NOTICE file
-# distributed with this work for additional information
-# regarding copyright ownership. The ASF licenses this file
-# to you under the Apache License, Version 2.0 (the
-# "License"); you may not use this file except in compliance
-# with the License. You may obtain a copy of the License at
-#
-# http://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing,
-# software distributed under the License is distributed on an
-# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
-# KIND, either express or implied. See the License for the
-# specific language governing permissions and limitations
-# under the License.
-"""cython3 namespace"""
diff --git a/python/tvm/_ffi/_cython/base.pxi b/python/tvm/_ffi/_cython/base.pxi
deleted file mode 100644
index 036a9c9559..0000000000
--- a/python/tvm/_ffi/_cython/base.pxi
+++ /dev/null
@@ -1,236 +0,0 @@
-# Licensed to the Apache Software Foundation (ASF) under one
-# or more contributor license agreements. See the NOTICE file
-# distributed with this work for additional information
-# regarding copyright ownership. The ASF licenses this file
-# to you under the Apache License, Version 2.0 (the
-# "License"); you may not use this file except in compliance
-# with the License. You may obtain a copy of the License at
-#
-# http://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing,
-# software distributed under the License is distributed on an
-# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
-# KIND, either express or implied. See the License for the
-# specific language governing permissions and limitations
-# under the License.
-
-from ..base import raise_last_ffi_error
-from libcpp cimport bool as bool_t
-from libcpp.vector cimport vector
-from cpython.version cimport PY_MAJOR_VERSION
-from cpython cimport pycapsule
-from libc.stdint cimport int32_t, int64_t, uint64_t, uint32_t, uint8_t, uint16_t
-import ctypes
-
-cdef enum TVMArgTypeCode:
- kInt = 0
- kUInt = 1
- kFloat = 2
- kTVMOpaqueHandle = 3
- kTVMNullptr = 4
- kTVMDataType = 5
- kDLDevice = 6
- kTVMDLTensorHandle = 7
- kTVMObjectHandle = 8
- kTVMModuleHandle = 9
- kTVMPackedFuncHandle = 10
- kTVMStr = 11
- kTVMBytes = 12
- kTVMNDArrayHandle = 13
- kTVMObjectRefArg = 14
- kTVMArgBool = 15
- kTVMExtBegin = 16
-
-cdef extern from "tvm/runtime/c_runtime_api.h":
- ctypedef struct DLDataType:
- uint8_t code
- uint8_t bits
- uint16_t lanes
-
- ctypedef struct DLDevice:
- int device_type
- int device_id
-
- ctypedef struct DLTensor:
- void* data
- DLDevice device
- int ndim
- DLDataType dtype
- int64_t* shape
- int64_t* strides
- uint64_t byte_offset
-
- ctypedef struct DLManagedTensor:
- DLTensor dl_tensor
- void* manager_ctx
- void (*deleter)(DLManagedTensor* self)
-
- ctypedef struct TVMValue:
- int64_t v_int64
- bool_t v_bool
- double v_float64
- void* v_handle
- const char* v_str
- DLDataType v_type
- DLDevice v_device
-
-ctypedef int64_t tvm_index_t
-ctypedef DLTensor* DLTensorHandle
-ctypedef void* TVMStreamHandle
-ctypedef void* TVMRetValueHandle
-ctypedef void* TVMPackedFuncHandle
-ctypedef void* ObjectHandle
-
-ctypedef struct TVMObject:
- uint32_t type_index_
- int32_t ref_counter_
- void (*deleter_)(TVMObject* self)
-
-
-ctypedef int (*TVMPackedCFunc)(
- TVMValue* args,
- int* type_codes,
- int num_args,
- TVMRetValueHandle ret,
- void* resource_handle)
-
-ctypedef void (*TVMPackedCFuncFinalizer)(void* resource_handle)
-
-# NOTE: All of TVM's C API function can be called without gil.
-# for API functions that can be run long(e.g. FuncCall)
-# we need to explicitly release the GIL as follows.
-#
-# cdef myfunc():
-# cdef int c_api_ret_code
-# with nogil:
-# c_api_ret_code = TVMAPIFunc(...)
-# CHECK_CALL(c_apt_ret_code)
-#
-# Explicitly releasing the GIL enables other python threads
-# to continue running while we are in TVMAPIFunc.
-# Not releasing GIL explicitly is OK(and perhaps desirable)
-# for short-running functions, as frequent unlocking also takes time,
-# the python interpreter will release GIL in a set period.
-#
-# We mark the possibly long running function as nogil below.
-cdef extern from "tvm/runtime/c_runtime_api.h":
- void TVMAPISetLastError(const char* msg)
- void TVMAPISetLastPythonError(void* py_object) except +
- const char *TVMGetLastError()
- int TVMFuncGetGlobal(const char* name,
- TVMPackedFuncHandle* out)
- int TVMFuncCall(TVMPackedFuncHandle func,
- TVMValue* arg_values,
- int* type_codes,
- int num_args,
- TVMValue* ret_val,
- int* ret_type_code) nogil
- int TVMFuncFree(TVMPackedFuncHandle func)
- int TVMCFuncSetReturn(TVMRetValueHandle ret,
- TVMValue* value,
- int* type_code,
- int num_ret)
- int TVMFuncCreateFromCFunc(TVMPackedCFunc func,
- void* resource_handle,
- TVMPackedCFuncFinalizer fin,
- TVMPackedFuncHandle *out)
- int TVMCbArgToReturn(TVMValue* value, int* code)
- int TVMArrayAlloc(tvm_index_t* shape,
- tvm_index_t ndim,
- DLDataType dtype,
- DLDevice dev,
- DLTensorHandle* out) nogil
- int TVMArrayFree(DLTensorHandle handle) nogil
- int TVMArrayCopyFromTo(DLTensorHandle src,
- DLTensorHandle to,
- TVMStreamHandle stream) nogil
- int TVMArrayFromDLPack(DLManagedTensor* arr_from,
- DLTensorHandle* out) nogil
- int TVMArrayToDLPack(DLTensorHandle arr_from,
- DLManagedTensor** out) nogil
- void TVMDLManagedTensorCallDeleter(DLManagedTensor* dltensor)
- int TVMObjectFree(ObjectHandle obj)
- int TVMObjectGetTypeIndex(ObjectHandle obj, unsigned* out_index)
-
-
-# the new FFI C API
-cdef extern from "tvm/ffi/c_api.h":
- ctypedef void* TVMFFIObjectHandle
-
- ctypedef struct TVMFFIByteArray:
- const char* data
- int64_t size
-
- int TVMFFIObjectGetTypeIndex(TVMFFIObjectHandle obj) nogil
- TVMFFIByteArray* TVMFFIBytesGetByteArrayPtr(TVMFFIObjectHandle obj) nogil
-
-
-cdef inline py_str(const char* x):
- return x.decode("utf-8")
-
-
-cdef inline c_str(pystr):
- """Create ctypes char * from a python string
- Parameters
- ----------
- string : string type
- python string
-
- Returns
- -------
- str : c_char_p
- A char pointer that can be passed to C API
- """
- return pystr.encode("utf-8")
-
-
-cdef inline int CHECK_CALL(int ret) except -2:
- """Check the return code of the C API function call"""
- # -2 brings exception
- if ret == -2:
- return -2
- if ret != 0:
- raise_last_ffi_error()
- return 0
-
-
-cdef inline object ctypes_handle(void* chandle):
- """Cast C handle to ctypes handle."""
- return ctypes.cast(<unsigned long long>chandle, ctypes.c_void_p)
-
-
-cdef inline void* c_handle(object handle):
- """Cast C types handle to c handle."""
- cdef unsigned long long v_ptr
- v_ptr = handle.value
- return <void*>(v_ptr)
-
-
-# python env API
-cdef extern from "Python.h":
- int PyErr_CheckSignals()
- void* PyGILState_Ensure()
- void PyGILState_Release(void*)
- void Py_IncRef(void*)
- void Py_DecRef(void*)
-
-cdef extern from "tvm/runtime/c_backend_api.h":
- int TVMBackendRegisterEnvCAPI(const char* name, void* ptr)
-
-cdef _init_env_api():
- # Initialize env api for signal handling
- # so backend can call tvm::runtime::EnvCheckSignals to check
- # signal when executing a long running function.
- #
- # Also registers the gil state release and ensure as PyErr_CheckSignals
- # function is called with gil released and we need to regrab the gil
- CHECK_CALL(TVMBackendRegisterEnvCAPI(c_str("PyErr_CheckSignals"), <void*>PyErr_CheckSignals))
- CHECK_CALL(TVMBackendRegisterEnvCAPI(c_str("PyGILState_Ensure"), <void*>PyGILState_Ensure))
- CHECK_CALL(TVMBackendRegisterEnvCAPI(c_str("PyGILState_Release"), <void*>PyGILState_Release))
- CHECK_CALL(TVMBackendRegisterEnvCAPI(c_str("PyGILState_Release"), <void*>PyGILState_Release))
- CHECK_CALL(TVMBackendRegisterEnvCAPI(c_str("Py_IncRef"), <void*>Py_IncRef))
- CHECK_CALL(TVMBackendRegisterEnvCAPI(c_str("Py_DecRef"), <void*>Py_DecRef))
-
-
-_init_env_api()
diff --git a/python/tvm/_ffi/_cython/core.pyx b/python/tvm/_ffi/_cython/core.pyx
deleted file mode 100644
index 730f8fc133..0000000000
--- a/python/tvm/_ffi/_cython/core.pyx
+++ /dev/null
@@ -1,21 +0,0 @@
-# Licensed to the Apache Software Foundation (ASF) under one
-# or more contributor license agreements. See the NOTICE file
-# distributed with this work for additional information
-# regarding copyright ownership. The ASF licenses this file
-# to you under the Apache License, Version 2.0 (the
-# "License"); you may not use this file except in compliance
-# with the License. You may obtain a copy of the License at
-#
-# http://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing,
-# software distributed under the License is distributed on an
-# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
-# KIND, either express or implied. See the License for the
-# specific language governing permissions and limitations
-# under the License.
-
-include "./base.pxi"
-include "./object.pxi"
-include "./packed_func.pxi"
-include "./ndarray.pxi"
diff --git a/python/tvm/_ffi/_cython/ndarray.pxi b/python/tvm/_ffi/_cython/ndarray.pxi
deleted file mode 100644
index f220e866d3..0000000000
--- a/python/tvm/_ffi/_cython/ndarray.pxi
+++ /dev/null
@@ -1,180 +0,0 @@
-# Licensed to the Apache Software Foundation (ASF) under one
-# or more contributor license agreements. See the NOTICE file
-# distributed with this work for additional information
-# regarding copyright ownership. The ASF licenses this file
-# to you under the Apache License, Version 2.0 (the
-# "License"); you may not use this file except in compliance
-# with the License. You may obtain a copy of the License at
-#
-# http://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing,
-# software distributed under the License is distributed on an
-# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
-# KIND, either express or implied. See the License for the
-# specific language governing permissions and limitations
-# under the License.
-
-from ..runtime_ctypes import TVMArrayHandle
-from cpython cimport PyCapsule_Destructor
-
-cdef const char* _c_str_dltensor = "dltensor"
-cdef const char* _c_str_used_dltensor = "used_dltensor"
-
-
-cdef void _c_dlpack_deleter(object pycaps):
- cdef DLManagedTensor* dltensor
- if pycapsule.PyCapsule_IsValid(pycaps, _c_str_dltensor):
- dltensor = <DLManagedTensor*>pycapsule.PyCapsule_GetPointer(pycaps, _c_str_dltensor)
- dltensor.deleter(dltensor)
-
-
-def _from_dlpack(object dltensor):
- cdef DLManagedTensor* ptr
- cdef DLTensorHandle chandle
- cdef int c_api_ret_code
- if pycapsule.PyCapsule_IsValid(dltensor, _c_str_dltensor):
- ptr = <DLManagedTensor*>pycapsule.PyCapsule_GetPointer(dltensor, _c_str_dltensor)
- with nogil:
- c_api_ret_code = TVMArrayFromDLPack(ptr, &chandle)
- CHECK_CALL(c_api_ret_code)
- # set name and destructor to be empty
- pycapsule.PyCapsule_SetDestructor(dltensor, NULL)
- pycapsule.PyCapsule_SetName(dltensor, _c_str_used_dltensor)
- return c_make_array(chandle, False, False)
- raise ValueError("Expect a dltensor field, pycapsule.PyCapsule can only be consumed once")
-
-
-cdef class NDArrayBase:
- cdef DLTensor* chandle
- cdef int c_is_view
-
- cdef inline _set_handle(self, handle):
- cdef unsigned long long ptr
- if handle is None:
- self.chandle = NULL
- else:
- ptr = ctypes.cast(handle, ctypes.c_void_p).value
- self.chandle = <DLTensor*>(ptr)
-
- property _tvm_handle:
- def __get__(self):
- return <unsigned long long>self.chandle
-
- property handle:
- def __get__(self):
- if self.chandle == NULL:
- return None
- else:
- return ctypes.cast(
- <unsigned long long>self.chandle, TVMArrayHandle)
-
- def __set__(self, value):
- self._set_handle(value)
-
- property is_view:
- def __get__(self):
- return self.c_is_view != 0
-
- @property
- def shape(self):
- """Shape of this array"""
- return tuple(self.chandle.shape[i] for i in range(self.chandle.ndim))
-
- def __init__(self, handle, is_view):
- self._set_handle(handle)
- self.c_is_view = is_view
-
- def __dealloc__(self):
- cdef int c_api_ret_code
- if self.c_is_view == 0:
- with nogil:
- c_api_ret_code = TVMArrayFree(self.chandle)
- CHECK_CALL(c_api_ret_code)
-
- def _copyto(self, target_nd):
- """Internal function that implements copy to target ndarray."""
- cdef int c_api_ret_code
- with nogil:
- c_api_ret_code = TVMArrayCopyFromTo(self.chandle, (<NDArrayBase>target_nd).chandle, NULL)
- CHECK_CALL(c_api_ret_code)
- return target_nd
-
- def to_dlpack(self):
- """Produce an array from a DLPack Tensor without copying memory
-
- Returns
- -------
- dlpack : DLPack tensor view of the array data
- """
- cdef DLManagedTensor* dltensor
- cdef int c_api_ret_code
- if self.c_is_view != 0:
- raise ValueError("to_dlpack do not work with memory views")
- with nogil:
- c_api_ret_code = TVMArrayToDLPack(self.chandle, &dltensor)
- CHECK_CALL(c_api_ret_code)
- return pycapsule.PyCapsule_New(dltensor, _c_str_dltensor, <PyCapsule_Destructor>_c_dlpack_deleter)
-
-
-# Import limited object-related function from C++ side to improve the speed
-# NOTE: can only use POD-C compatible object in FFI.
-cdef extern from "tvm/runtime/ndarray.h" namespace "tvm::runtime":
- cdef void* TVMArrayHandleToObjectHandle(DLTensorHandle handle)
-
-
-cdef c_make_array(void* chandle, is_view, is_container):
- global _TVM_ND_CLS
-
- if is_container:
- tindex = (
- <TVMObject*>TVMArrayHandleToObjectHandle(<DLTensorHandle>chandle)).type_index_
- if tindex < len(_TVM_ND_CLS):
- cls = _TVM_ND_CLS[tindex]
- if cls is not None:
- ret = cls.__new__(cls)
- else:
- ret = _CLASS_NDARRAY.__new__(_CLASS_NDARRAY)
- else:
- ret = _CLASS_NDARRAY.__new__(_CLASS_NDARRAY)
- (<NDArrayBase>ret).chandle = <DLTensor*>chandle
- (<NDArrayBase>ret).c_is_view = <int>is_view
- return ret
- else:
- ret = _CLASS_NDARRAY.__new__(_CLASS_NDARRAY)
- (<NDArrayBase>ret).chandle = <DLTensor*>chandle
- (<NDArrayBase>ret).c_is_view = <int>is_view
- return ret
-
-
-cdef _TVM_COMPATS = ()
-
-cdef _TVM_EXT_RET = {}
-
-def _reg_extension(cls, fcreate):
- global _TVM_COMPATS
- _TVM_COMPATS += (cls,)
- if fcreate:
- _TVM_EXT_RET[cls._tvm_tcode] = fcreate
-
-cdef list _TVM_ND_CLS = []
-
-cdef _register_ndarray(int index, object cls):
- """register object class"""
- global _TVM_ND_CLS
- while len(_TVM_ND_CLS) <= index:
- _TVM_ND_CLS.append(None)
-
- _TVM_ND_CLS[index] = cls
-
-
-def _make_array(handle, is_view, is_container):
- cdef unsigned long long ptr
- ptr = ctypes.cast(handle, ctypes.c_void_p).value
- return c_make_array(<void*>ptr, is_view, is_container)
-
-cdef object _CLASS_NDARRAY = None
-
-def _set_class_ndarray(cls):
- global _CLASS_NDARRAY
- _CLASS_NDARRAY = cls
diff --git a/python/tvm/_ffi/_cython/object.pxi b/python/tvm/_ffi/_cython/object.pxi
deleted file mode 100644
index e45a13365c..0000000000
--- a/python/tvm/_ffi/_cython/object.pxi
+++ /dev/null
@@ -1,165 +0,0 @@
-# Licensed to the Apache Software Foundation (ASF) under one
-# or more contributor license agreements. See the NOTICE file
-# distributed with this work for additional information
-# regarding copyright ownership. The ASF licenses this file
-# to you under the Apache License, Version 2.0 (the
-# "License"); you may not use this file except in compliance
-# with the License. You may obtain a copy of the License at
-#
-# http://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing,
-# software distributed under the License is distributed on an
-# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
-# KIND, either express or implied. See the License for the
-# specific language governing permissions and limitations
-# under the License.
-
-"""Maps object type index to its constructor"""
-cdef list OBJECT_TYPE = []
-"""Maps object type to its type index"""
-cdef dict OBJECT_INDEX = {}
-
-def _register_object(int index, object cls):
- """register object class"""
- if issubclass(cls, NDArrayBase):
- _register_ndarray(index, cls)
- return
-
- global OBJECT_TYPE
- while len(OBJECT_TYPE) <= index:
- OBJECT_TYPE.append(None)
- OBJECT_TYPE[index] = cls
- OBJECT_INDEX[cls] = index
-
-def _get_object_type_index(object cls):
- """get the type index of object class"""
- return OBJECT_INDEX.get(cls)
-
-cdef inline object make_ret_object(void* chandle):
- global OBJECT_TYPE
- global _CLASS_OBJECT
- cdef unsigned tindex
- cdef object cls
- cdef object handle
- object_type = OBJECT_TYPE
- handle = ctypes_handle(chandle)
- tindex = TVMFFIObjectGetTypeIndex(chandle)
-
- if tindex < len(OBJECT_TYPE):
- cls = OBJECT_TYPE[tindex]
- if cls is not None:
- if issubclass(cls, PyNativeObject):
- obj = _CLASS_OBJECT.__new__(_CLASS_OBJECT)
- (<ObjectBase>obj).chandle = chandle
- return cls.__from_tvm_object__(cls, obj)
- obj = cls.__new__(cls)
- else:
- obj = _CLASS_OBJECT.__new__(_CLASS_OBJECT)
- else:
- obj = _CLASS_OBJECT.__new__(_CLASS_OBJECT)
-
- (<ObjectBase>obj).chandle = chandle
-
- # Handle return values that must be converted from the TVM object
- # to a python native object. This should be used in cases where
- # subclassing the python native object is forbidden. For example,
- # `runtime.BoxBool` cannot be a subclass of `bool`, as `bool` does
- # not allow any subclasses.
- # if hasattr(obj, '__into_pynative_object__'):
- # return obj.__into_pynative_object__)
-
- return obj
- # return obj.__into_pynative_object__()
-
-
-class PyNativeObject:
- """Base class of all TVM objects that also subclass python's builtin types."""
- __slots__ = []
-
- def __init_tvm_object_by_constructor__(self, fconstructor, *args):
- """Initialize the internal tvm_object by calling constructor function.
-
- Parameters
- ----------
- fconstructor : Function
- Constructor function.
-
- args: list of objects
- The arguments to the constructor
-
- Note
- ----
- We have a special calling convention to call constructor functions.
- So the return object is directly set into the object
- """
- obj = _CLASS_OBJECT.__new__(_CLASS_OBJECT)
- obj.__init_handle_by_constructor__(fconstructor, *args)
- self.__tvm_object__ = obj
-
-
-cdef class ObjectBase:
- cdef void* chandle
-
- cdef inline _set_handle(self, handle):
- cdef unsigned long long ptr
- if handle is None:
- self.chandle = NULL
- else:
- ptr = handle.value
- self.chandle = <void*>(ptr)
-
- property handle:
- def __get__(self):
- return ctypes_handle(self.chandle)
-
- def __set__(self, value):
- self._set_handle(value)
-
- def __dealloc__(self):
- CHECK_CALL(TVMObjectFree(self.chandle))
-
- def __init_handle_by_constructor__(self, fconstructor, *args):
- """Initialize the handle by calling constructor function.
-
- Parameters
- ----------
- fconstructor : Function
- Constructor function.
-
- args: list of objects
- The arguments to the constructor
-
- Note
- ----
- We have a special calling convention to call constructor functions.
- So the return handle is directly set into the Node object
- instead of creating a new Node.
- """
- # avoid error raised during construction.
- self.chandle = NULL
- cdef void* chandle
- ConstructorCall(
- (<PackedFuncBase>fconstructor).chandle, args, &chandle)
- self.chandle = chandle
-
- def same_as(self, other):
- """Check object identity.
-
- Parameters
- ----------
- other : object
- The other object to compare against.
-
- Returns
- -------
- result : bool
- The comparison result.
- """
- if not isinstance(other, ObjectBase):
- return False
- return self.chandle == (<ObjectBase>other).chandle
-
-def StringGetPyString(obj):
- cdef TVMFFIByteArray* bytes = TVMFFIBytesGetByteArrayPtr((<ObjectBase>obj).chandle)
- return py_str(bytes.data)
diff --git a/python/tvm/_ffi/_cython/packed_func.pxi b/python/tvm/_ffi/_cython/packed_func.pxi
deleted file mode 100644
index 19af447a11..0000000000
--- a/python/tvm/_ffi/_cython/packed_func.pxi
+++ /dev/null
@@ -1,376 +0,0 @@
-# Licensed to the Apache Software Foundation (ASF) under one
-# or more contributor license agreements. See the NOTICE file
-# distributed with this work for additional information
-# regarding copyright ownership. The ASF licenses this file
-# to you under the Apache License, Version 2.0 (the
-# "License"); you may not use this file except in compliance
-# with the License. You may obtain a copy of the License at
-#
-# http://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing,
-# software distributed under the License is distributed on an
-# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
-# KIND, either express or implied. See the License for the
-# specific language governing permissions and limitations
-# under the License.
-
-import ctypes
-import traceback
-from cpython cimport Py_INCREF, Py_DECREF, PyGILState_Ensure, PyGILState_Release
-from numbers import Number, Integral
-from ..base import string_types, py2cerror
-from ..runtime_ctypes import DataType, Device, TVMByteArray, ObjectRValueRef
-
-
-cdef void tvm_callback_finalize(void* fhandle) with gil:
- local_pyfunc = <object>(fhandle)
- Py_DECREF(local_pyfunc)
-
-cdef int tvm_callback(TVMValue* args,
- int* type_codes,
- int num_args,
- TVMRetValueHandle ret,
- void* fhandle) with gil:
- cdef list pyargs
- cdef TVMValue value
- cdef int tcode
- local_pyfunc = <object>(fhandle)
- pyargs = []
- for i in range(num_args):
- value = args[i]
- tcode = type_codes[i]
- if (tcode == kTVMObjectHandle or
- tcode == kTVMPackedFuncHandle or
- tcode == kTVMModuleHandle or
- tcode == kTVMNDArrayHandle or
- tcode == kTVMObjectRefArg or
- tcode >= kTVMExtBegin):
- CHECK_CALL(TVMCbArgToReturn(&value, &tcode))
-
- if tcode != kTVMDLTensorHandle:
- pyargs.append(make_ret(value, tcode))
- else:
- pyargs.append(c_make_array(value.v_handle, True, False))
- try:
- rv = local_pyfunc(*pyargs)
- except Exception as err:
- msg = traceback.format_exc()
- msg = py2cerror(msg)
- TVMAPISetLastPythonError(<void*>err)
-
- return -1
- if rv is not None:
- if isinstance(rv, tuple):
- raise ValueError("PackedFunction can only support one return value")
- temp_args = []
- make_arg(rv, &value, &tcode, temp_args)
- CHECK_CALL(TVMCFuncSetReturn(ret, &value, &tcode, 1))
- return 0
-
-
-cdef object make_packed_func(TVMPackedFuncHandle chandle, int is_global):
- obj = _CLASS_PACKED_FUNC.__new__(_CLASS_PACKED_FUNC)
- (<PackedFuncBase>obj).chandle = chandle
- (<PackedFuncBase>obj).is_global = is_global
- return obj
-
-
-def convert_to_tvm_func(object pyfunc):
- """Convert a python function to TVM function
-
- Parameters
- ----------
- pyfunc : python function
- The python function to be converted.
-
- Returns
- -------
- tvmfunc: tvm.Function
- The converted tvm function.
- """
- cdef TVMPackedFuncHandle chandle
- Py_INCREF(pyfunc)
- CHECK_CALL(TVMFuncCreateFromCFunc(tvm_callback,
- <void*>(pyfunc),
- tvm_callback_finalize,
- &chandle))
- return make_packed_func(chandle, False)
-
-
-cdef inline int make_arg(object arg,
- TVMValue* value,
- int* tcode,
- list temp_args) except -1:
- """Pack arguments into c args tvm call accept"""
- cdef unsigned long long ptr
- if isinstance(arg, ObjectBase):
- value[0].v_handle = (<ObjectBase>arg).chandle
- tcode[0] = kTVMObjectHandle
- elif isinstance(arg, NDArrayBase):
- value[0].v_handle = (<NDArrayBase>arg).chandle
- tcode[0] = (kTVMNDArrayHandle if
- not (<NDArrayBase>arg).c_is_view else kTVMDLTensorHandle)
- elif isinstance(arg, PyNativeObject):
- value[0].v_handle = (<ObjectBase>(arg.__tvm_object__)).chandle
- tcode[0] = kTVMObjectHandle
- elif isinstance(arg, _TVM_COMPATS):
- ptr = arg._tvm_handle
- value[0].v_handle = (<void*>ptr)
- tcode[0] = arg.__class__._tvm_tcode
- elif isinstance(arg, bool):
- # A python `bool` is a subclass of `int`, so this check
- # must occur before `Integral`.
- value[0].v_int64 = arg
- tcode[0] = kTVMArgBool
- elif isinstance(arg, Integral):
- value[0].v_int64 = arg
- tcode[0] = kInt
- elif isinstance(arg, float):
- value[0].v_float64 = arg
- tcode[0] = kFloat
- elif isinstance(arg, str):
- tstr = c_str(arg)
- value[0].v_str = tstr
- tcode[0] = kTVMStr
- temp_args.append(tstr)
- elif arg is None:
- value[0].v_handle = NULL
- tcode[0] = kTVMNullptr
- elif isinstance(arg, Number):
- value[0].v_float64 = arg
- tcode[0] = kFloat
- elif isinstance(arg, DataType):
- tstr = c_str(str(arg))
- value[0].v_str = tstr
- tcode[0] = kTVMStr
- temp_args.append(tstr)
- elif isinstance(arg, Device):
- value[0].v_device = (<DLDevice*>(
- <unsigned long long>ctypes.addressof(arg)))[0]
- tcode[0] = kDLDevice
- elif isinstance(arg, (bytes, bytearray)):
- # from_buffer only taeks in bytearray.
- if isinstance(arg, bytes):
- byte_arr = bytearray(arg)
- temp_args.append(byte_arr)
- arg = byte_arr
-
- arr = TVMByteArray()
- arr.data = ctypes.cast(
- (ctypes.c_byte * len(arg)).from_buffer(arg),
- ctypes.POINTER(ctypes.c_byte))
- arr.size = len(arg)
- value[0].v_handle = <void*>(
- <unsigned long long>ctypes.addressof(arr))
- tcode[0] = kTVMBytes
- temp_args.append(arr)
- elif isinstance(arg, string_types):
- tstr = c_str(arg)
- value[0].v_str = tstr
- tcode[0] = kTVMStr
- temp_args.append(tstr)
- elif isinstance(arg, (list, tuple, dict, _CLASS_OBJECT_GENERIC)):
- arg = _FUNC_CONVERT_TO_OBJECT(arg)
- value[0].v_handle = (<ObjectBase>arg).chandle
- tcode[0] = kTVMObjectHandle
- temp_args.append(arg)
- elif isinstance(arg, _CLASS_MODULE):
- value[0].v_handle = c_handle(arg.handle)
- tcode[0] = kTVMModuleHandle
- elif isinstance(arg, PackedFuncBase):
- value[0].v_handle = (<PackedFuncBase>arg).chandle
- tcode[0] = kTVMPackedFuncHandle
- elif isinstance(arg, ctypes.c_void_p):
- value[0].v_handle = c_handle(arg)
- tcode[0] = kTVMOpaqueHandle
- elif isinstance(arg, ObjectRValueRef):
- value[0].v_handle = &((<ObjectBase>(arg.obj)).chandle)
- tcode[0] = kTVMObjectRefArg
- elif callable(arg):
- arg = convert_to_tvm_func(arg)
- value[0].v_handle = (<PackedFuncBase>arg).chandle
- tcode[0] = kTVMPackedFuncHandle
- temp_args.append(arg)
- else:
- raise TypeError("Don't know how to handle type %s" % type(arg))
- return 0
-
-
-cdef inline bytearray make_ret_bytes(void* chandle):
- handle = ctypes_handle(chandle)
- arr = ctypes.cast(handle, ctypes.POINTER(TVMByteArray))[0]
- size = arr.size
- res = bytearray(size)
- rptr = (ctypes.c_byte * size).from_buffer(res)
- if not ctypes.memmove(rptr, arr.data, size):
- raise RuntimeError('memmove failed')
- return res
-
-
-cdef inline object make_ret(TVMValue value, int tcode):
- """convert result to return value."""
- if tcode == kTVMObjectHandle:
- return make_ret_object(value.v_handle)
- elif tcode == kTVMNullptr:
- return None
- elif tcode == kTVMArgBool:
- return bool(value.v_int64)
- elif tcode == kInt:
- return value.v_int64
- elif tcode == kFloat:
- return value.v_float64
- elif tcode == kTVMNDArrayHandle:
- return c_make_array(value.v_handle, False, True)
- elif tcode == kTVMStr:
- return py_str(value.v_str)
- elif tcode == kTVMBytes:
- return make_ret_bytes(value.v_handle)
- elif tcode == kTVMOpaqueHandle:
- return ctypes_handle(value.v_handle)
- elif tcode == kDLDevice:
- return Device(value.v_device.device_type, value.v_device.device_id)
- elif tcode == kTVMModuleHandle:
- return _CLASS_MODULE(ctypes_handle(value.v_handle))
- elif tcode == kTVMPackedFuncHandle:
- return make_packed_func(value.v_handle, False)
- elif tcode in _TVM_EXT_RET:
- return _TVM_EXT_RET[tcode](ctypes_handle(value.v_handle))
-
- raise ValueError("Unhandled type code %d" % tcode)
-
-
-cdef inline int FuncCall3(void* chandle,
- tuple args,
- int nargs,
- TVMValue* ret_val,
- int* ret_tcode) except -1:
- cdef TVMValue[3] values
- cdef int[3] tcodes
- nargs = len(args)
- temp_args = []
- for i in range(nargs):
- make_arg(args[i], &values[i], &tcodes[i], temp_args)
-
- with nogil:
- c_api_ret_code = TVMFuncCall(chandle, &values[0], &tcodes[0],
- nargs, ret_val, ret_tcode)
-
- CHECK_CALL(c_api_ret_code)
- return 0
-
-cdef inline int FuncCall(void* chandle,
- tuple args,
- TVMValue* ret_val,
- int* ret_tcode) except -1:
- cdef int nargs
- cdef int c_api_ret_code
- nargs = len(args)
- if nargs <= 3:
- FuncCall3(chandle, args, nargs, ret_val, ret_tcode)
- return 0
-
- cdef vector[TVMValue] values
- cdef vector[int] tcodes
- values.resize(max(nargs, 1))
- tcodes.resize(max(nargs, 1))
- temp_args = []
- for i in range(nargs):
- make_arg(args[i], &values[i], &tcodes[i], temp_args)
-
- with nogil:
- c_api_ret_code = TVMFuncCall(chandle, &values[0], &tcodes[0],
- nargs, ret_val, ret_tcode)
- CHECK_CALL(c_api_ret_code)
- return 0
-
-
-cdef inline int ConstructorCall(void* constructor_handle,
- tuple args,
- void** handle) except -1:
- """Call contructor of a handle function"""
- cdef TVMValue ret_val
- cdef int ret_tcode
- FuncCall(constructor_handle, args, &ret_val, &ret_tcode)
- handle[0] = ret_val.v_handle
- return 0
-
-
-cdef class PackedFuncBase:
- cdef TVMPackedFuncHandle chandle
- cdef int is_global
-
- cdef inline _set_handle(self, handle):
- if handle is None:
- self.chandle = NULL
- else:
- self.chandle = c_handle(handle)
-
- property is_global:
- def __get__(self):
- return self.c_is_global != 0
-
- def __set__(self, value):
- self.c_is_global = value
-
- property handle:
- def __get__(self):
- if self.chandle == NULL:
- return None
- else:
- return ctypes.cast(<unsigned long long>self.chandle, ctypes.c_void_p)
- def __set__(self, value):
- self._set_handle(value)
-
- def __init__(self, handle, is_global):
- self._set_handle(handle)
- self.c_is_global = is_global
-
- def __dealloc__(self):
- if self.is_global == 0:
- CHECK_CALL(TVMFuncFree(self.chandle))
-
- def __call__(self, *args):
- cdef TVMValue ret_val
- cdef int ret_tcode
- ret_tcode = kTVMNullptr
- FuncCall(self.chandle, args, &ret_val, &ret_tcode)
- return make_ret(ret_val, ret_tcode)
-
-
-def _get_global_func(name, allow_missing):
- cdef TVMPackedFuncHandle chandle
- CHECK_CALL(TVMFuncGetGlobal(c_str(name), &chandle))
- if chandle != NULL:
- return make_packed_func(chandle, True)
-
- if allow_missing:
- return None
-
- raise ValueError("Cannot find global function %s" % name)
-
-
-_CLASS_PACKED_FUNC = None
-_CLASS_MODULE = None
-_CLASS_OBJECT = None
-_CLASS_OBJECT_GENERIC = None
-_FUNC_CONVERT_TO_OBJECT = None
-
-def _set_class_module(module_class):
- """Initialize the module."""
- global _CLASS_MODULE
- _CLASS_MODULE = module_class
-
-def _set_class_packed_func(func_class):
- global _CLASS_PACKED_FUNC
- _CLASS_PACKED_FUNC = func_class
-
-def _set_class_object(obj_class):
- global _CLASS_OBJECT
- _CLASS_OBJECT = obj_class
-
-def _set_class_object_generic(object_generic_class, func_convert_to_object):
- global _CLASS_OBJECT_GENERIC
- global _FUNC_CONVERT_TO_OBJECT
- _CLASS_OBJECT_GENERIC = object_generic_class
- _FUNC_CONVERT_TO_OBJECT = func_convert_to_object
(tvm) 02/06: [FFI] Simplify unpack traits
Posted by tq...@apache.org.
This is an automated email from the ASF dual-hosted git repository.
tqchen pushed a commit to branch refactor-s3
in repository https://gitbox.apache.org/repos/asf/tvm.git
commit 0ddf789068eae441a2cdee9abec35be641d17d76
Author: tqchen <ti...@gmail.com>
AuthorDate: Sun May 4 20:49:32 2025 -0400
[FFI] Simplify unpack traits
---
ffi/include/tvm/ffi/container/variant.h | 4 +-
ffi/include/tvm/ffi/function.h | 27 +++++----
ffi/include/tvm/ffi/function_details.h | 93 ++++-------------------------
ffi/include/tvm/ffi/reflection/reflection.h | 9 +--
include/tvm/ir/env_func.h | 11 +++-
include/tvm/runtime/packed_func.h | 19 +++---
src/runtime/metal/metal_module.mm | 2 +-
src/tir/ir/expr.cc | 55 ++++++++---------
src/tir/schedule/instruction_traits.h | 9 +--
9 files changed, 90 insertions(+), 139 deletions(-)
diff --git a/ffi/include/tvm/ffi/container/variant.h b/ffi/include/tvm/ffi/container/variant.h
index 1455a5b34a..c8b58ba49e 100644
--- a/ffi/include/tvm/ffi/container/variant.h
+++ b/ffi/include/tvm/ffi/container/variant.h
@@ -94,12 +94,12 @@ class Variant {
template <typename T, typename = enable_if_variant_contains_t<T>>
TVM_FFI_INLINE T get() const& {
- return data_.cast<T>();
+ return data_.template cast<T>();
}
template <typename T, typename = enable_if_variant_contains_t<T>>
TVM_FFI_INLINE T get() && {
- return std::move(data_).cast<T>();
+ return std::move(data_).template cast<T>();
}
TVM_FFI_INLINE std::string GetTypeKey() const { return data_.GetTypeKey(); }
diff --git a/ffi/include/tvm/ffi/function.h b/ffi/include/tvm/ffi/function.h
index 913237ec04..e4b67e4a76 100644
--- a/ffi/include/tvm/ffi/function.h
+++ b/ffi/include/tvm/ffi/function.h
@@ -453,8 +453,8 @@ class Function : public ObjectRef {
static Function FromUnpacked(TCallable callable) {
using FuncInfo = details::FunctionInfo<TCallable>;
auto call_packed = [callable](const AnyView* args, int32_t num_args, Any* rv) mutable -> void {
- details::unpack_call<typename FuncInfo::RetType, FuncInfo::num_args>(nullptr, callable, args,
- num_args, rv);
+ details::unpack_call<typename FuncInfo::RetType>(
+ std::make_index_sequence<FuncInfo::num_args>{}, nullptr, callable, args, num_args, rv);
};
return FromPackedInternal(call_packed);
}
@@ -469,8 +469,8 @@ class Function : public ObjectRef {
using FuncInfo = details::FunctionInfo<TCallable>;
auto call_packed = [callable, name](const AnyView* args, int32_t num_args,
Any* rv) mutable -> void {
- details::unpack_call<typename FuncInfo::RetType, FuncInfo::num_args>(&name, callable, args,
- num_args, rv);
+ details::unpack_call<typename FuncInfo::RetType>(
+ std::make_index_sequence<FuncInfo::num_args>{}, &name, callable, args, num_args, rv);
};
return FromPackedInternal(call_packed);
}
@@ -674,7 +674,16 @@ class TypedFunction<R(Args...)> {
* \returns The return value.
*/
TVM_FFI_INLINE R operator()(Args... args) const {
- return details::typed_packed_call_dispatcher<R>::run(packed_, std::forward<Args>(args)...);
+ if constexpr (std::is_same_v<R, void>) {
+ packed_(std::forward<Args>(args)...);
+ } else {
+ Any res = packed_(std::forward<Args>(args)...);
+ if constexpr (std::is_same_v<R, Any>) {
+ return res;
+ } else {
+ return std::move(res).cast<R>();
+ }
+ }
}
/*!
* \brief convert to PackedFunc
@@ -850,8 +859,6 @@ class Function::Registry {
return *this;
}
- operator details::EmptyStruct() const { return details::EmptyStruct(); }
-
protected:
/*!
* \brief set the body of the function to be f
@@ -875,8 +882,8 @@ inline int32_t TypeKeyToIndex(const char* type_key) {
return type_index;
}
-#define TVM_FFI_REG_VAR_DEF \
- static inline TVM_FFI_ATTRIBUTE_UNUSED ::tvm::ffi::details::EmptyStruct __mk_##TVMFFI
+#define TVM_FFI_FUNC_REG_VAR_DEF \
+ static inline TVM_FFI_ATTRIBUTE_UNUSED ::tvm::ffi::Function::Registry& __##TVMFFIFuncReg
/*!
* \brief Register a function globally.
@@ -888,7 +895,7 @@ inline int32_t TypeKeyToIndex(const char* type_key) {
* \endcode
*/
#define TVM_FFI_REGISTER_GLOBAL(OpName) \
- TVM_FFI_STR_CONCAT(TVM_FFI_REG_VAR_DEF, __COUNTER__) = ::tvm::ffi::Function::Registry(OpName)
+ TVM_FFI_STR_CONCAT(TVM_FFI_FUNC_REG_VAR_DEF, __COUNTER__) = ::tvm::ffi::Function::Registry(OpName)
} // namespace ffi
} // namespace tvm
#endif // TVM_FFI_FUNCTION_H_
diff --git a/ffi/include/tvm/ffi/function_details.h b/ffi/include/tvm/ffi/function_details.h
index 34e166428e..f47a253a58 100644
--- a/ffi/include/tvm/ffi/function_details.h
+++ b/ffi/include/tvm/ffi/function_details.h
@@ -36,14 +36,6 @@ namespace tvm {
namespace ffi {
namespace details {
-/*!
- * \brief Empty struct, used to reduce unused variable in global static initialization.
- */
-struct EmptyStruct {
- /*! \brief at least one byte to ensure address */
- uint8_t pad;
-};
-
template <typename ArgType>
struct Arg2Str {
template <size_t i>
@@ -167,46 +159,10 @@ class ArgValueWithContext {
FGetFuncSignature f_sig_;
};
-template <typename R, int nleft, int index, typename F>
-struct unpack_call_dispatcher {
- template <typename... Args>
- TVM_FFI_INLINE static void run(const std::string* optional_name, FGetFuncSignature f_sig,
- const F& f, const AnyView* args, int32_t num_args, Any* rv,
- Args&&... unpacked_args) {
- // construct a movable argument value
- // which allows potential move of argument to the input of F.
- unpack_call_dispatcher<R, nleft - 1, index + 1, F>::run(
- optional_name, f_sig, f, args, num_args, rv, std::forward<Args>(unpacked_args)...,
- ArgValueWithContext(args, index, optional_name, f_sig));
- }
-};
-
-template <typename R, int index, typename F>
-struct unpack_call_dispatcher<R, 0, index, F> {
- template <typename... Args>
- TVM_FFI_INLINE static void run(const std::string*, FGetFuncSignature, const F& f, const AnyView*,
- int32_t, Any* rv, Args&&... unpacked_args) {
- using RetType = decltype(f(std::forward<Args>(unpacked_args)...));
- if constexpr (std::is_same_v<RetType, R>) {
- *rv = f(std::forward<Args>(unpacked_args)...);
- } else {
- *rv = R(f(std::forward<Args>(unpacked_args)...));
- }
- }
-};
-
-template <int index, typename F>
-struct unpack_call_dispatcher<void, 0, index, F> {
- template <typename... Args>
- TVM_FFI_INLINE static void run(const std::string*, FGetFuncSignature, const F& f, const AnyView*,
- int32_t, Any*, Args&&... unpacked_args) {
- f(std::forward<Args>(unpacked_args)...);
- }
-};
-
-template <typename R, int nargs, typename F>
-TVM_FFI_INLINE void unpack_call(const std::string* optional_name, const F& f, const AnyView* args,
- int32_t num_args, Any* rv) {
+template <typename R, std::size_t... Is, typename F>
+TVM_FFI_INLINE void unpack_call(std::index_sequence<Is...>, const std::string* optional_name,
+ const F& f, [[maybe_unused]] const AnyView* args,
+ [[maybe_unused]] int32_t num_args, [[maybe_unused]] Any* rv) {
using FuncInfo = FunctionInfo<F>;
FGetFuncSignature f_sig = FuncInfo::Sig;
@@ -214,47 +170,20 @@ TVM_FFI_INLINE void unpack_call(const std::string* optional_name, const F& f, co
#ifndef _MSC_VER
static_assert(FuncInfo::unpacked_supported, "The function signature do not support unpacked");
#endif
-
+ constexpr size_t nargs = sizeof...(Is);
if (nargs != num_args) {
TVM_FFI_THROW(TypeError) << "Mismatched number of arguments when calling: `"
<< (optional_name == nullptr ? "" : *optional_name)
<< (f_sig == nullptr ? "" : (*f_sig)()) << "`. Expected " << nargs
<< " but got " << num_args << " arguments";
}
- unpack_call_dispatcher<R, nargs, 0, F>::run(optional_name, f_sig, f, args, num_args, rv);
-}
-
-template <typename FType>
-struct unpack_call_by_signature {};
-
-template <typename R, typename... Args>
-struct unpack_call_by_signature<R(Args...)> {
- template <typename F>
- TVM_FFI_INLINE static void run(const F& f, const AnyView* args, int32_t num_args, Any* rv) {
- unpack_call<R, sizeof...(Args)>(nullptr, f, args, num_args, rv);
+ // use index sequence to do recursive-less unpacking
+ if constexpr (std::is_same_v<R, void>) {
+ f(ArgValueWithContext(args, Is, optional_name, f_sig)...);
+ } else {
+ *rv = R(f(ArgValueWithContext(args, Is, optional_name, f_sig)...));
}
-};
-
-template <typename R>
-struct typed_packed_call_dispatcher {
- template <typename F, typename... Args>
- TVM_FFI_INLINE static R run(const F& f, Args&&... args) {
- Any res = f(std::forward<Args>(args)...);
- if constexpr (std::is_same_v<R, Any>) {
- return res;
- } else {
- return std::move(res).cast<R>();
- }
- }
-};
-
-template <>
-struct typed_packed_call_dispatcher<void> {
- template <typename F, typename... Args>
- TVM_FFI_INLINE static void run(const F& f, Args&&... args) {
- f(std::forward<Args>(args)...);
- }
-};
+}
/*!
* \brief Move the safe call raised error to the caller
diff --git a/ffi/include/tvm/ffi/reflection/reflection.h b/ffi/include/tvm/ffi/reflection/reflection.h
index b1088c1eda..8ce2d22ddd 100644
--- a/ffi/include/tvm/ffi/reflection/reflection.h
+++ b/ffi/include/tvm/ffi/reflection/reflection.h
@@ -66,8 +66,6 @@ class ReflectionDef {
return *this;
}
- operator details::EmptyStruct() const { return details::EmptyStruct(); }
-
private:
template <typename Class, typename T>
void RegisterField(const char* name, T Class::*field_ptr, bool readonly) {
@@ -138,11 +136,14 @@ class ReflectionFieldGetter {
const TVMFFIFieldInfo* field_info_;
};
+#define TVM_FFI_REFLECTION_REG_VAR_DEF \
+ static inline TVM_FFI_ATTRIBUTE_UNUSED ::tvm::ffi::details::ReflectionDef& __TVMFFIReflectionReg
+
/*!
* helper macro to define a reflection definition for an object
*/
-#define TVM_FFI_REFLECTION_DEF(TypeName) \
- TVM_FFI_STR_CONCAT(TVM_FFI_REG_VAR_DEF, __COUNTER__) = \
+#define TVM_FFI_REFLECTION_DEF(TypeName) \
+ TVM_FFI_STR_CONCAT(TVM_FFI_REFLECTION_REG_VAR_DEF, __COUNTER__) = \
::tvm::ffi::details::ReflectionDef(TypeName::_GetOrAllocRuntimeTypeIndex())
} // namespace details
} // namespace ffi
diff --git a/include/tvm/ir/env_func.h b/include/tvm/ir/env_func.h
index c83711756b..c44e102cca 100644
--- a/include/tvm/ir/env_func.h
+++ b/include/tvm/ir/env_func.h
@@ -139,7 +139,16 @@ class TypedEnvFunc<R(Args...)> : public ObjectRef {
R operator()(Args... args) const {
const EnvFuncNode* n = operator->();
ICHECK(n != nullptr);
- return ffi::details::typed_packed_call_dispatcher<R>::run(n->func, std::forward<Args>(args)...);
+ if constexpr (std::is_same_v<R, void>) {
+ n->func(std::forward<Args>(args)...);
+ } else {
+ Any res = n->func(std::forward<Args>(args)...);
+ if constexpr (std::is_same_v<R, Any>) {
+ return res;
+ } else {
+ return std::move(res).cast<R>();
+ }
+ }
}
/*! \brief specify container node */
using ContainerType = EnvFuncNode;
diff --git a/include/tvm/runtime/packed_func.h b/include/tvm/runtime/packed_func.h
index 228ce731e0..3609987d55 100644
--- a/include/tvm/runtime/packed_func.h
+++ b/include/tvm/runtime/packed_func.h
@@ -343,7 +343,8 @@ struct ModuleVTableEntryHelper<R (T::*)(Args...) const> {
using MemFnType = R (T::*)(Args...) const;
static TVM_ALWAYS_INLINE void Call(TVMRetValue* rv, T* self, MemFnType f, TVMArgs args) {
auto wrapped = [self, f](Args... args) -> R { return (self->*f)(std::forward<Args>(args)...); };
- ffi::details::unpack_call<R, sizeof...(Args)>(nullptr, wrapped, args.data(), args.size(), rv);
+ ffi::details::unpack_call<R>(std::make_index_sequence<sizeof...(Args)>{}, nullptr, wrapped,
+ args.data(), args.size(), rv);
}
};
@@ -352,7 +353,8 @@ struct ModuleVTableEntryHelper<R (T::*)(Args...)> {
using MemFnType = R (T::*)(Args...);
static TVM_ALWAYS_INLINE void Call(TVMRetValue* rv, T* self, MemFnType f, TVMArgs args) {
auto wrapped = [self, f](Args... args) -> R { return (self->*f)(std::forward<Args>(args)...); };
- ffi::details::unpack_call<R, sizeof...(Args)>(nullptr, wrapped, args.data(), args.size(), rv);
+ ffi::details::unpack_call<R>(std::make_index_sequence<sizeof...(Args)>{}, nullptr, wrapped,
+ args.data(), args.size(), rv);
}
};
@@ -361,8 +363,8 @@ struct ModuleVTableEntryHelper<void (T::*)(Args...) const> {
using MemFnType = void (T::*)(Args...) const;
static TVM_ALWAYS_INLINE void Call(TVMRetValue* rv, T* self, MemFnType f, TVMArgs args) {
auto wrapped = [self, f](Args... args) -> void { (self->*f)(std::forward<Args>(args)...); };
- ffi::details::unpack_call<void, sizeof...(Args)>(nullptr, wrapped, args.data(), args.size(),
- rv);
+ ffi::details::unpack_call<void>(std::make_index_sequence<sizeof...(Args)>{}, nullptr, wrapped,
+ args.data(), args.size(), rv);
}
};
@@ -371,8 +373,8 @@ struct ModuleVTableEntryHelper<void (T::*)(Args...)> {
using MemFnType = void (T::*)(Args...);
static TVM_ALWAYS_INLINE void Call(TVMRetValue* rv, T* self, MemFnType f, TVMArgs args) {
auto wrapped = [self, f](Args... args) -> void { (self->*f)(std::forward<Args>(args)...); };
- ffi::details::unpack_call<void, sizeof...(Args)>(nullptr, wrapped, args.data(), args.size(),
- rv);
+ ffi::details::unpack_call<void>(std::make_index_sequence<sizeof...(Args)>{}, nullptr, wrapped,
+ args.data(), args.size(), rv);
}
};
} // namespace details
@@ -446,8 +448,9 @@ struct ModuleVTableEntryHelper<void (T::*)(Args...)> {
TVM_FFI_SAFE_CALL_BEGIN(); \
using FuncInfo = ::tvm::ffi::details::FunctionInfo<decltype(Function)>; \
static std::string name = #ExportName; \
- ::tvm::ffi::details::unpack_call<typename FuncInfo::RetType, FuncInfo::num_args>( \
- &name, Function, reinterpret_cast<const ::tvm::ffi::AnyView*>(args), num_args, \
+ ::tvm::ffi::details::unpack_call<typename FuncInfo::RetType>( \
+ std::make_index_sequence<FuncInfo::num_args>{}, &name, Function, \
+ reinterpret_cast<const ::tvm::ffi::AnyView*>(args), num_args, \
reinterpret_cast<::tvm::ffi::Any*>(result)); \
TVM_FFI_SAFE_CALL_END(); \
} \
diff --git a/src/runtime/metal/metal_module.mm b/src/runtime/metal/metal_module.mm
index b338274231..a9a7faaefa 100644
--- a/src/runtime/metal/metal_module.mm
+++ b/src/runtime/metal/metal_module.mm
@@ -211,7 +211,7 @@ class MetalWrappedFunc {
id<MTLComputeCommandEncoder> encoder = [cb computeCommandEncoder];
[encoder setComputePipelineState:scache_[device_id]];
for (size_t i = 0; i < num_buffer_args_; ++i) {
- void* buf = args[static_cast<int>(i)];
+ void* buf = args[static_cast<int>(i)].cast<void*>();
[encoder setBuffer:(id<MTLBuffer>)(buf) offset:0 atIndex:i];
}
if (num_pack_args_ != 0) {
diff --git a/src/tir/ir/expr.cc b/src/tir/ir/expr.cc
index 491488ba30..bec6c04085 100644
--- a/src/tir/ir/expr.cc
+++ b/src/tir/ir/expr.cc
@@ -558,36 +558,37 @@ Call::Call(DataType dtype, RelaxExpr op, Array<PrimExpr> args, Span span) {
}
TVM_REGISTER_GLOBAL("tir.Call")
- .set_body_typed([](Optional<DataType> dtype, RelaxExpr op,
- Array<Variant<runtime::String, DLDataType, IterVar, BufferRegion, PrimExpr>> args,
- Span span) {
- Array<PrimExpr> prim_expr_args;
- for (const auto& it : args) {
- if (auto opt_str = it.as<String>()) {
- prim_expr_args.push_back(StringImm(opt_str.value()));
- } else if (auto opt_dtype = it.as<DLDataType>()) {
- prim_expr_args.push_back(StringImm(ffi::DLDataTypeToString(opt_dtype.value())));
- } else if (const auto* iter_var = it.as<IterVarNode>()) {
- prim_expr_args.push_back(iter_var->var);
- } else if (const auto* br = it.as<BufferRegionNode>()) {
- Array<PrimExpr> indices;
- for (Range r : br->region) {
- if (is_one(r->extent)) {
- indices.push_back(r->min);
- } else if (r->extent.as<IntImmNode>()) {
- indices.push_back(tir::Ramp(r->min, make_const(r->min->dtype, 1), r->extent));
+ .set_body_typed(
+ [](Optional<DataType> dtype, RelaxExpr op,
+ Array<Variant<runtime::String, DLDataType, IterVar, BufferRegion, PrimExpr>> args,
+ Span span) {
+ Array<PrimExpr> prim_expr_args;
+ for (const auto& it : args) {
+ if (auto opt_str = it.as<String>()) {
+ prim_expr_args.push_back(StringImm(opt_str.value()));
+ } else if (auto opt_dtype = it.as<DLDataType>()) {
+ prim_expr_args.push_back(StringImm(ffi::DLDataTypeToString(opt_dtype.value())));
+ } else if (const auto* iter_var = it.as<IterVarNode>()) {
+ prim_expr_args.push_back(iter_var->var);
+ } else if (const auto* br = it.as<BufferRegionNode>()) {
+ Array<PrimExpr> indices;
+ for (Range r : br->region) {
+ if (is_one(r->extent)) {
+ indices.push_back(r->min);
+ } else if (r->extent.as<IntImmNode>()) {
+ indices.push_back(tir::Ramp(r->min, make_const(r->min->dtype, 1), r->extent));
+ } else {
+ LOG(FATAL) << "ValueError: Cannot convert to BufferLoad: "
+ << GetRef<BufferRegion>(br);
+ }
+ }
+ prim_expr_args.push_back(BufferLoad(br->buffer, indices));
} else {
- LOG(FATAL) << "ValueError: Cannot convert to BufferLoad: "
- << GetRef<BufferRegion>(br);
+ prim_expr_args.push_back(Downcast<PrimExpr>(it));
}
}
- prim_expr_args.push_back(BufferLoad(br->buffer, indices));
- } else {
- prim_expr_args.push_back(Downcast<PrimExpr>(it));
- }
- }
- return Call(dtype.value_or(DataType::Void()), op, prim_expr_args, span);
- });
+ return Call(dtype.value_or(DataType::Void()), op, prim_expr_args, span);
+ });
TVM_REGISTER_NODE_TYPE(CallNode);
diff --git a/src/tir/schedule/instruction_traits.h b/src/tir/schedule/instruction_traits.h
index d292bd3dc7..29165750e8 100644
--- a/src/tir/schedule/instruction_traits.h
+++ b/src/tir/schedule/instruction_traits.h
@@ -319,8 +319,9 @@ Array<Any> UnpackedInstTraits<TTraits>::ApplyToSchedule(const Schedule& sch,
PackedFunc pf([](const TVMArgs& args, TVMRetValue* rv) -> void {
constexpr size_t kNumArgs = details::NumArgs<method_type>;
ICHECK_EQ(args.size(), kNumArgs);
- ffi::details::unpack_call<return_type, kNumArgs>(nullptr, TTraits::UnpackedApplyToSchedule,
- args.data(), args.size(), rv);
+ ffi::details::unpack_call<return_type>(std::make_index_sequence<kNumArgs>{}, nullptr,
+ TTraits::UnpackedApplyToSchedule, args.data(),
+ args.size(), rv);
});
ffi::Any rv;
pf.CallPacked(ffi::PackedArgs(packed_args, kNumArgs), &rv);
@@ -349,8 +350,8 @@ String UnpackedInstTraits<TTraits>::AsPython(const Array<Any>& inputs, const Arr
PackedFunc pf([](const TVMArgs& args, TVMRetValue* rv) -> void {
constexpr size_t kNumArgs = details::NumArgs<method_type>;
ICHECK_EQ(args.size(), kNumArgs);
- ffi::details::unpack_call<return_type, kNumArgs>(nullptr, TTraits::UnpackedAsPython,
- args.data(), args.size(), rv);
+ ffi::details::unpack_call<return_type>(std::make_index_sequence<kNumArgs>{}, nullptr,
+ TTraits::UnpackedAsPython, args.data(), args.size(), rv);
});
ffi::Any rv;
pf.CallPacked(ffi::PackedArgs(packed_args, kNumArgs), &rv);