This is an automated email from the ASF dual-hosted git repository.
guan404ming pushed a commit to branch main
in repository https://gitbox.apache.org/repos/asf/tvm.git
The following commit(s) were added to refs/heads/main by this push:
new bbfdab79d9 [CI] Repair Python test cleanup regressions (#19955)
bbfdab79d9 is described below
commit bbfdab79d998cfc6af22b0be510c543501d447b7
Author: Tianqi Chen <[email protected]>
AuthorDate: Mon Jul 6 16:29:52 2026 +0800
[CI] Repair Python test cleanup regressions (#19955)
## Summary
- Keep the Python test launcher close to plain `pytest -n auto`, move
nightly tests under `tests/nightly/python`, remove obsolete launchers
and collection bookkeeping, and partition CPU/GPU jobs with explicit
`gpu` marker expressions.
- Repair exact-pointer regressions at their owning boundaries: packed
raw-string ABI values, CUDA/Metal matrix intrinsic pointers, internal TE
extern offsets, MetaSchedule scalar annotations, localized
auto-tensorization scope matching, and typed DLTensor fixture fields.
- Preserve typed workspace calls in TIR and cast pointer-returning
external calls in CodeGenC, covered by a plain-TIRx 1024-byte global
workspace that is compiled as C++.
- Finish phasing out value-bearing Relax `R.Prim` annotations by
requiring an explicit dtype, removing obsolete value-based contracts,
and expressing the DISCO rank-dependent slices as explicit scalar
`call_tir` inputs.
- Gate the distributed callback on the optional DISCO runtime, NCCL, and
at least two GPUs so capability-limited jobs skip instead of failing.
- Remove the non-demonstrating pointer probe, use direct TVMScript
comparison for packed strings, and remove the four designated legacy
testing modules.
The seven repaired CPU categories cover packed raw strings (7 failures),
CUDA/Metal matrix access-pointer types (7), internal TE extern offsets
(1), a typed DLTensor fixture (1), MetaSchedule scalar annotations (1),
CodeGenC workspace return casts (12), and localized auto-tensorization
storage-scope matching (19).
## Validation
- Base: `ded6ad8dd212869c881efb5590f8a33fc972728e`
- Head: `a7277e86dbcfe0638c8c252d36760859c4ab4297`
- All 35 locally available original failing node IDs pass across the
focused runs.
- The full focused TE, TIR builtin-lowering, and CodeGenC files pass: 61
tests.
- The complete touched Relax/TVMScript set plus
PlanAndUpdateBufferAllocationLocation passes with 784 passed, 20
skipped, and 1 expected failure.
- The DISCO callback collects and skips when its runtime or two-GPU
environment is unavailable.
- Six direct mapping tests, twelve tensor-core sketches, and the dp4a
sketch pass unchanged.
- The compiler rebuild, branch-wide pre-commit hooks, and full-range
whitespace checks pass.
- The 13 broad CBLAS/TFLite nodes remain dependency-gated; their owning
TE and generated-C regressions compile.
No merge is included in this change.
---
ci/jenkins/generated/gpu_jenkinsfile.groovy | 4 -
ci/jenkins/templates/gpu_jenkinsfile.groovy.j2 | 4 -
conftest.py | 32 ---
docs/conf.py | 2 +-
docs/contribute/code_guide.rst | 5 +-
docs/contribute/pull_request.rst | 2 +-
docs/contribute/testing.rst | 4 +-
pyproject.toml | 5 +-
python/tvm/relax/script/parser/entry.py | 22 +-
python/tvm/s_tir/tensor_intrin/cuda.py | 4 +-
python/tvm/s_tir/tensor_intrin/metal.py | 4 +-
python/tvm/testing/plugin.py | 104 ----------
python/tvm/testing/utils.py | 124 +++--------
src/s_tir/meta_schedule/mutator/mutate_parallel.cc | 2 +-
src/s_tir/schedule/ir_comparator.cc | 18 +-
src/target/source/codegen_c.cc | 9 +
src/te/operation/create_primfunc.cc | 11 +-
src/tirx/op/op.cc | 6 +
src/tirx/transform/lower_tvm_builtin.cc | 1 +
.../python}/test_nnapi/__init__.py | 0
.../python}/test_nnapi/conftest.py | 0
.../python}/test_nnapi/infrastructure.py | 0
.../test_nnapi/test_from_exported_to_cuda.py | 0
.../python}/test_nnapi/test_network.py | 0
.../python}/test_nnapi/test_ops.py | 0
tests/python/codegen/test_target_codegen_c_host.py | 18 ++
tests/python/conftest.py | 21 +-
tests/python/disco/test_callback.py | 69 +++++--
tests/python/disco/test_ccl.py | 3 +
tests/python/disco/test_custom_allreduce.py | 6 +-
tests/python/relax/test_analysis_type_analysis.py | 50 +----
.../relax/test_backend_transform_shape_lower.py | 81 --------
tests/python/relax/test_bind_params.py | 22 +-
tests/python/relax/test_bind_symbolic_vars.py | 67 ------
tests/python/relax/test_blockbuilder_core.py | 8 +-
tests/python/relax/test_dataflow_rewriter.py | 149 --------------
tests/python/relax/test_op_binary.py | 32 ---
.../relax/test_transform_compute_prim_value.py | 31 ---
.../relax/test_transform_lazy_transform_params.py | 145 -------------
.../relax/test_transform_lift_transform_params.py | 4 +-
.../test_transform_remove_unused_parameters.py | 21 +-
tests/python/relax/test_tvmscript_parser.py | 8 +
tests/python/relax/test_utils.py | 62 +++---
tests/python/relax/test_vm_build.py | 94 ---------
.../request_hook => python}/request_hook.py | 4 +-
...sform_plan_update_buffer_allocation_location.py | 8 +-
tests/python/te/test_te_create_primfunc.py | 37 ++++
tests/python/testing/test_testing.py | 116 -----------
.../testing/test_tvm_testing_before_after.py | 147 -------------
tests/python/testing/test_tvm_testing_features.py | 192 -----------------
.../python/testing/test_type_annotation_checker.py | 227 ---------------------
.../test_tir_transform_lower_tvm_builtin.py | 30 +++
tests/python/tirx/conftest.py | 13 +-
tests/python/tvmscript/test_tvmscript_roundtrip.py | 10 +-
tests/scripts/ci.py | 5 -
tests/scripts/task_clear_pytest.sh | 22 --
tests/scripts/task_python_integration.sh | 32 ---
tests/scripts/task_python_integration_gpuonly.sh | 27 ---
tests/scripts/task_python_unittest.sh | 50 +----
tests/scripts/task_python_unittest_gpuonly.sh | 11 +-
60 files changed, 339 insertions(+), 1846 deletions(-)
diff --git a/ci/jenkins/generated/gpu_jenkinsfile.groovy
b/ci/jenkins/generated/gpu_jenkinsfile.groovy
index 8624efbea8..e7ce14c4ba 100644
--- a/ci/jenkins/generated/gpu_jenkinsfile.groovy
+++ b/ci/jenkins/generated/gpu_jenkinsfile.groovy
@@ -530,10 +530,6 @@ def run_unittest_GPU(node_type) {
script: "${docker_run} ${ci_gpu}
./tests/scripts/task_python_unittest_gpuonly.sh",
label: 'Run Python GPU unit tests',
)
- sh (
- script: "${docker_run} ${ci_gpu}
./tests/scripts/task_python_integration_gpuonly.sh",
- label: 'Run Python GPU integration tests',
- )
})
}
}
diff --git a/ci/jenkins/templates/gpu_jenkinsfile.groovy.j2
b/ci/jenkins/templates/gpu_jenkinsfile.groovy.j2
index b05e84c7e7..4553fbefa2 100644
--- a/ci/jenkins/templates/gpu_jenkinsfile.groovy.j2
+++ b/ci/jenkins/templates/gpu_jenkinsfile.groovy.j2
@@ -53,10 +53,6 @@
script: "${docker_run} ${ci_gpu}
./tests/scripts/task_python_unittest_gpuonly.sh",
label: 'Run Python GPU unit tests',
)
- sh (
- script: "${docker_run} ${ci_gpu}
./tests/scripts/task_python_integration_gpuonly.sh",
- label: 'Run Python GPU integration tests',
- )
{% endcall %}
diff --git a/conftest.py b/conftest.py
deleted file mode 100644
index dd294ea548..0000000000
--- a/conftest.py
+++ /dev/null
@@ -1,32 +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 os
-import sys
-from pathlib import Path
-
-pytest_plugins = ["tvm.testing.plugin"]
-IS_IN_CI = os.getenv("CI", "") == "true"
-REPO_ROOT = Path(__file__).resolve().parent
-
-
-def pytest_sessionstart():
- if IS_IN_CI:
- hook_script_dir = REPO_ROOT / "tests" / "scripts" / "request_hook"
- sys.path.append(str(hook_script_dir))
- import request_hook # pylint: disable=import-outside-toplevel
-
- request_hook.init()
diff --git a/docs/conf.py b/docs/conf.py
index 24f43e1eb6..c226de1c40 100644
--- a/docs/conf.py
+++ b/docs/conf.py
@@ -215,7 +215,7 @@ def rst2md(text, gallery_conf, target_dir, heading_levels,
real_func):
def install_request_hook(gallery_conf, fname):
- testing.utils.install_request_hook(depth=3)
+ testing.utils.install_request_hook(tvm_path.resolve() / "tests" / "python"
/ "request_hook.py")
INSTALL_TVM_DEV = """\
diff --git a/docs/contribute/code_guide.rst b/docs/contribute/code_guide.rst
index 1622af120c..5136b9d4d6 100644
--- a/docs/contribute/code_guide.rst
+++ b/docs/contribute/code_guide.rst
@@ -127,7 +127,8 @@ Python Code Styles
Writing Python Tests
--------------------
-We use `pytest <https://docs.pytest.org/en/stable/>`_ for all python testing.
``tests/python`` contains all the tests.
+We use `pytest <https://docs.pytest.org/en/stable/>`_ for all Python testing.
Regular tests live
+under ``tests/python``, while environment-specific nightly tests live under
``tests/nightly/python``.
See :doc:`testing` for details on running tests, target parametrization,
and the target-specific marks used by CI.
@@ -162,7 +163,7 @@ server can go down or be slow), so try to avoid using the
network at all during
this isn't a reasonable proposition (e.g. the docs tutorials which need to
download models).
New network downloads are rejected by the CI `request hook
-<https://github.com/apache/tvm/blob/main/tests/scripts/request_hook/request_hook.py>`_.
Prefer
+<https://github.com/apache/tvm/blob/main/tests/python/request_hook.py>`_.
Prefer
checked-in fixtures or generated data. If a download is unavoidable, arrange a
stable
project-managed mirror with the maintainers and add an explicit ``URL_MAP``
entry.
diff --git a/docs/contribute/pull_request.rst b/docs/contribute/pull_request.rst
index a6c41b755f..9150838e76 100644
--- a/docs/contribute/pull_request.rst
+++ b/docs/contribute/pull_request.rst
@@ -251,7 +251,7 @@ If you want to run all tests:
# build tvm (see install-from-source for CMake build instructions)
cd build && cmake .. && cmake --build . --parallel $(nproc) && cd ..
- ./tests/scripts/task_python_unittest.sh
+ python -m pytest -vvs -n auto tests/python
If you want to run a single test:
diff --git a/docs/contribute/testing.rst b/docs/contribute/testing.rst
index 74a308945a..f30fd7d03b 100644
--- a/docs/contribute/testing.rst
+++ b/docs/contribute/testing.rst
@@ -299,8 +299,8 @@ in which stages.
- Which tests run
- The ``Unit Test`` and ``Integration Test`` stages of the Jenkinsfile
- determine how ``pytest`` is called. Each task starts by unpacking a
+ The ``Unit Test`` stage of the Jenkinsfile determines how ``pytest``
+ is called. Each task starts by unpacking a
compiled library that was previous compiled in the ``BUILD`` stage,
then runs a test script
(e.g. ``tests/scripts/task_python_unittest.sh``). These scripts set
diff --git a/pyproject.toml b/pyproject.toml
index 2c38d1f9d8..a4f80182cf 100644
--- a/pyproject.toml
+++ b/pyproject.toml
@@ -167,10 +167,7 @@ addopts = "-v --tb=short"
python_files = ["test_*.py", "*_test.py"]
python_classes = ["Test*"]
python_functions = ["test_*"]
-markers = [
- "gpu: Mark a test as requiring a GPU",
- "xdist_group(name): Run related tests in one pytest-xdist load group",
-]
+markers = ["gpu: Mark a test as requiring a GPU"]
[tool.ruff]
include = [
diff --git a/python/tvm/relax/script/parser/entry.py
b/python/tvm/relax/script/parser/entry.py
index 1812a1102e..20605840f6 100644
--- a/python/tvm/relax/script/parser/entry.py
+++ b/python/tvm/relax/script/parser/entry.py
@@ -449,32 +449,21 @@ def Shape(values: list[Expr] | None = None, ndim: int =
-1) -> ShapeProxy:
class PrimProxy(TypeProxy):
- dtype: str | None
+ dtype: str
"""The type of TIR-representable values.
Parameters
----------
- dtype : Optional[str]
+ dtype : str
The data type.
"""
def __init__(
self,
- dtype: str | None = None,
- value: int | float | str | Expr | None = None,
+ dtype: str,
) -> None:
- if dtype is None:
- if tvm.ir.is_prim_expr(value):
- dtype = str(value.ty)
- elif isinstance(value, float):
- dtype = "float32"
- elif value is not None:
- dtype = "int64"
- else:
- raise TypeError("R.Prim missing required argument 'dtype'")
-
self.dtype = dtype
def get_symbolic_vars(self) -> set[str]:
@@ -485,10 +474,9 @@ class PrimProxy(TypeProxy):
def Prim(
- dtype: str | None = None,
- value: int | float | str | Expr | None = None,
+ dtype: str,
) -> PrimProxy:
- return PrimProxy(dtype, value)
+ return PrimProxy(dtype)
############################ R.match_cast #############################
diff --git a/python/tvm/s_tir/tensor_intrin/cuda.py
b/python/tvm/s_tir/tensor_intrin/cuda.py
index 8c28fd7ddf..701ac5e4a3 100644
--- a/python/tvm/s_tir/tensor_intrin/cuda.py
+++ b/python/tvm/s_tir/tensor_intrin/cuda.py
@@ -889,7 +889,7 @@ def get_wmma_load_intrin(
n_dim,
k_dim,
get_wmma_fragment_index(C, d1, frag_m, frag_n),
- A.access_ptr("r"),
+ A.access_ptr("r", ptr_type=dtype),
s1,
layout,
dtype="void",
@@ -1016,7 +1016,7 @@ def get_wmma_store_intrin(
n_dim,
k_dim,
get_wmma_fragment_index(A, d1, m_dim, n_dim),
- C.access_ptr("w"),
+ C.access_ptr("w", ptr_type=dtype),
s1,
"row_major",
dtype="void",
diff --git a/python/tvm/s_tir/tensor_intrin/metal.py
b/python/tvm/s_tir/tensor_intrin/metal.py
index 132906b67d..5b19793ddf 100644
--- a/python/tvm/s_tir/tensor_intrin/metal.py
+++ b/python/tvm/s_tir/tensor_intrin/metal.py
@@ -125,7 +125,7 @@ def get_simdgroup_load_intrin(
T.metal.simdgroup_load(
C.data,
index=get_simdgroup_index(C, d1, col, row),
- ptr=A.access_ptr("r"),
+ ptr=A.access_ptr("r", ptr_type=dtype),
stride=s1,
col=col,
row=row,
@@ -182,7 +182,7 @@ def get_simdgroup_store_intrin(
T.metal.simdgroup_store(
A.data,
index=get_simdgroup_index(A, s1, col, row),
- ptr=C.access_ptr("w"),
+ ptr=C.access_ptr("w", ptr_type=dtype),
stride=d1,
col=col,
row=row,
diff --git a/python/tvm/testing/plugin.py b/python/tvm/testing/plugin.py
deleted file mode 100644
index 2869664842..0000000000
--- a/python/tvm/testing/plugin.py
+++ /dev/null
@@ -1,104 +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.
-# pylint: disable=unused-argument
-
-"""Pytest plugin for using tvm testing extensions.
-
-TVM provides utilities for testing across all supported targets, and
-to more easily parametrize across many inputs. For more information
-on usage of these features, see documentation in the tvm.testing
-module.
-
-These are enabled by default in all pytests provided by tvm, but may
-be useful externally for one-off testing. To enable, add the
-following line to the test script, or to the conftest.py in the same
-directory as the test scripts.
-
- pytest_plugins = ['tvm.testing.plugin']
-
-"""
-
-import _pytest
-
-
-def pytest_collection_modifyitems(config, items):
- """Called after all tests are chosen, currently used for bookkeeping."""
- # pylint: disable=unused-argument
- _count_num_fixture_uses(items)
- _remove_global_fixture_definitions(items)
- _sort_tests(items)
-
-
-def _count_num_fixture_uses(items):
- # Helper function, counts the number of tests that use each cached
- # fixture. Should be called from pytest_collection_modifyitems().
- for item in items:
- is_skipped = item.get_closest_marker("skip") or any(
- mark.args[0] for mark in item.iter_markers("skipif")
- )
- if is_skipped:
- continue
-
- for fixturedefs in item._fixtureinfo.name2fixturedefs.values():
- # Only increment the active fixturedef, in a name has been
overridden.
- fixturedef = fixturedefs[-1]
- if hasattr(fixturedef.func, "num_tests_use_this_fixture"):
- fixturedef.func.num_tests_use_this_fixture[0] += 1
-
-
-def _remove_global_fixture_definitions(items):
- # Helper function, removes fixture definitions from the global
- # variables of the modules they were defined in. This is intended
- # to improve readability of error messages by giving a NameError
- # if a test function accesses a pytest fixture but doesn't include
- # it as an argument. Should be called from
- # pytest_collection_modifyitems().
-
- modules = set(item.module for item in items)
-
- for module in modules:
- for name in dir(module):
- obj = getattr(module, name)
- if hasattr(obj, "_pytestfixturefunction") and isinstance(
- obj._pytestfixturefunction,
_pytest.fixtures.FixtureFunctionMarker
- ):
- delattr(module, name)
-
-
-def _sort_tests(items):
- """Sort tests by file/function.
-
- By default, pytest will sort tests to maximize the re-use of
- fixtures. However, this assumes that all fixtures have an equal
- cost to generate, and no caches outside of those managed by
- pytest. A tvm.testing.parameter is effectively free, while
- reference data for testing may be quite large. Since most of the
- TVM fixtures are specific to a python function, sort the test
- ordering by python function, so that
- tvm.testing.utils._fixture_cache can be cleared sooner rather than
- later.
-
- Should be called from pytest_collection_modifyitems.
-
- """
-
- def sort_key(item):
- filename, lineno, test_name = item.location
- test_name = test_name.split("[")[0]
- return filename, lineno, test_name
-
- items.sort(key=sort_key)
diff --git a/python/tvm/testing/utils.py b/python/tvm/testing/utils.py
index 6b1c02b01e..f41f06dc07 100644
--- a/python/tvm/testing/utils.py
+++ b/python/tvm/testing/utils.py
@@ -24,7 +24,7 @@ Organization
This file contains functions expected to be called directly by a user
while writing unit tests. Integrations with the pytest framework
-are in plugin.py.
+for TVM's own test suite are in ``tests/python/conftest.py``.
Testing Markers
***************
@@ -73,6 +73,7 @@ import logging
import os
import pickle
import platform
+import runpy
import sys
import time
from pathlib import Path
@@ -91,6 +92,7 @@ from tvm.support import nvcc
SKIP_SLOW_TESTS = os.getenv("SKIP_SLOW_TESTS", "").lower() in {"true", "1",
"yes"}
IS_IN_CI = os.getenv("CI", "") == "true"
+_REQUEST_HOOK_INITIALIZERS = {}
skip_if_wheel_test = pytest.mark.skipif(
os.getenv("WHEEL_TEST", "").lower() in {"true", "1", "yes"},
@@ -679,10 +681,11 @@ def fixture(func=None, *, cache_return_value=False):
If the setup is expensive to perform, then the
cache_return_value=True argument can be passed to cache the setup.
The fixture function will be run only once (or once per parameter,
- if used with tvm.testing.parameter), and the same return value
- will be passed to all tests that use it. If the environment
- variable TVM_TEST_DISABLE_CACHE is set to a non-zero value, it
- will disable this feature and no caching will be performed.
+ if used with tvm.testing.parameter). The cached setup value is
+ retained for the lifetime of the test process, and each test receives
+ an independent copy. If the environment variable TVM_TEST_DISABLE_CACHE
+ is set to a non-zero value, it will disable this feature and no caching
+ will be performed.
Example
-------
@@ -719,15 +722,10 @@ def fixture(func=None, *, cache_return_value=False):
force_disable_cache = bool(int(os.environ.get("TVM_TEST_DISABLE_CACHE",
"0")))
cache_return_value = cache_return_value and not force_disable_cache
- # Deliberately at function scope, so that caching can track how
- # many times the fixture has been used. If used, the cache gets
- # cleared after the fixture is no longer needed.
- scope = "function"
-
def wraps(func):
if cache_return_value:
func = _fixture_cache(func)
- func = pytest.fixture(func, scope=scope)
+ func = pytest.fixture(func, scope="function")
return func
if func is None:
@@ -791,13 +789,6 @@ class _DeepCopyAllowedClasses(dict):
def _fixture_cache(func):
cache = {}
- # Can't use += on a bound method's property. Therefore, this is a
- # list rather than a variable so that it can be accessed from the
- # pytest_collection_modifyitems().
- num_tests_use_this_fixture = [0]
-
- num_times_fixture_used = 0
-
# Using functools.lru_cache would require the function arguments
# to be hashable, which wouldn't allow caching fixtures that
# depend on numpy arrays. For example, a fixture that takes a
@@ -821,41 +812,21 @@ def _fixture_cache(func):
@functools.wraps(func)
def wrapper(*args, **kwargs):
- if num_tests_use_this_fixture[0] == 0:
- raise RuntimeError(
- "Fixture use count is 0. "
- "This can occur if tvm.testing.plugin isn't registered. "
- "If using outside of the TVM test directory, "
- "please add `pytest_plugins = ['tvm.testing.plugin']` to your
conftest.py"
- )
+ cache_key = get_cache_key(*args, **kwargs)
try:
- cache_key = get_cache_key(*args, **kwargs)
-
- try:
- cached_value = cache[cache_key]
- except KeyError:
- cached_value = cache[cache_key] = func(*args, **kwargs)
-
- yield copy.deepcopy(
- cached_value,
- # allowed_class_list should be a list of classes that
- # are safe to copy using copy.deepcopy, but do not
- # implement __deepcopy__, __reduce__, or
- # __reduce_ex__.
- _DeepCopyAllowedClasses(allowed_class_list=[]),
- )
-
- finally:
- # Clear the cache once all tests that use a particular fixture
- # have completed.
- nonlocal num_times_fixture_used
- num_times_fixture_used += 1
- if num_times_fixture_used >= num_tests_use_this_fixture[0]:
- cache.clear()
-
- # Set in the pytest_collection_modifyitems(), by _count_num_fixture_uses
- wrapper.num_tests_use_this_fixture = num_tests_use_this_fixture
+ cached_value = cache[cache_key]
+ except KeyError:
+ cached_value = cache[cache_key] = func(*args, **kwargs)
+
+ return copy.deepcopy(
+ cached_value,
+ # allowed_class_list should be a list of classes that
+ # are safe to copy using copy.deepcopy, but do not
+ # implement __deepcopy__, __reduce__, or
+ # __reduce_ex__.
+ _DeepCopyAllowedClasses(allowed_class_list=[]),
+ )
return wrapper
@@ -893,49 +864,22 @@ def is_ampere_or_newer():
return major >= 8 and minor != 9
-def install_request_hook(depth: int) -> None:
- """Add a wrapper around urllib.request for CI tests"""
+def install_request_hook(hook_script: Path) -> None:
+ """Add a wrapper around urllib.request for CI tests."""
if not IS_IN_CI:
return
- #
https://sphinx-gallery.github.io/stable/faq.html#why-is-file-not-defined-what-can-i-use
- base = None
- msg = ""
+ hook_script = Path(hook_script).resolve()
+ if not hook_script.is_file():
+ raise RuntimeError(f"Request hook {hook_script} does not exist")
+
+ # Load the exact hook file without exposing the test root as an import
path.
+ # Cache its initializer because Sphinx invokes this once per gallery
example.
try:
- base = __file__
- msg += f"found file {__file__}\n"
- except NameError:
- msg += "no file\n"
-
- if base is None:
- hook_script_dir = Path.cwd().resolve()
- msg += "used path.cwd()\n"
- else:
- hook_script_dir = Path(base).resolve().parent
- msg += "used base()\n"
-
- msg += f"using depth {depth}\n"
- if depth <= 0:
- raise ValueError(f"depth less than 1 not supported, found: {depth}")
-
- # Go up the parent directories
- while depth > 0:
- msg += f"[depth={depth}] dir={hook_script_dir}\n"
- hook_script_dir = hook_script_dir.parent
- depth -= 1
-
- # Ensure the specified dir is valid
- hook_script_dir = hook_script_dir / "tests" / "scripts" / "request_hook"
- if not hook_script_dir.exists():
- raise RuntimeError(f"Directory {hook_script_dir} does not
exist:\n{msg}")
-
- # Import the hook and start it up (it's not included here directly to avoid
- # keeping a database of URLs inside the tvm Python package
- sys.path.append(str(hook_script_dir))
- # This import is intentionally delayed since it should only happen in CI
- import request_hook # pylint: disable=import-outside-toplevel
-
- request_hook.init()
+ init = _REQUEST_HOOK_INITIALIZERS[hook_script]
+ except KeyError:
+ init = _REQUEST_HOOK_INITIALIZERS[hook_script] =
runpy.run_path(str(hook_script))["init"]
+ init()
def strtobool(val):
diff --git a/src/s_tir/meta_schedule/mutator/mutate_parallel.cc
b/src/s_tir/meta_schedule/mutator/mutate_parallel.cc
index a0e3639ff9..29fcc72bfd 100644
--- a/src/s_tir/meta_schedule/mutator/mutate_parallel.cc
+++ b/src/s_tir/meta_schedule/mutator/mutate_parallel.cc
@@ -248,7 +248,7 @@ bool FindParallelDecision(const Trace& trace, TRandState*
rand_state,
get_sblock_insts.at(ann_inst->inputs[0].as_or_throw<s_tir::SBlockRV>().get());
TVM_FFI_ICHECK_EQ(get_sblock_inst->attrs.size(), 2);
candidate->inst = ffi::GetRef<Instruction>(ann_inst);
- candidate->parallel_extent =
ann_inst->inputs[1].as_or_throw<IntImm>()->value;
+ candidate->parallel_extent = ann_inst->inputs[1].cast<IntImm>()->value;
candidate->block_name = get_sblock_inst->attrs[0].as_or_throw<ffi::String>();
candidate->func_name = get_sblock_inst->attrs[1].as_or_throw<ffi::String>();
return true;
diff --git a/src/s_tir/schedule/ir_comparator.cc
b/src/s_tir/schedule/ir_comparator.cc
index 399fa1721d..e20bfcbf21 100644
--- a/src/s_tir/schedule/ir_comparator.cc
+++ b/src/s_tir/schedule/ir_comparator.cc
@@ -739,8 +739,22 @@ bool AutoTensorizeComparator::CompareBuffer(const Buffer&
lhs, const Buffer& rhs
if (it != rhs_buffer_map_.end()) {
equal = (*it).second.same_as(lhs);
} else {
- // Remap both buffer itself and buffer data, skip buffer shape and scope
- equal = DefEqual(lhs->data, rhs->data) && lhs->dtype == rhs->dtype;
+ // Remap both buffer itself and buffer data, skipping buffer shape and
storage scope. Auto
+ // tensorization inserts the cache stages that move workload buffers into
an intrinsic's
+ // required scope, while the pointer element type must still agree.
+ auto data_it = equal_map_.find(lhs->data);
+ if (data_it != equal_map_.end()) {
+ equal = data_it->second.same_as(rhs->data);
+ } else {
+ const auto* lhs_ptr = lhs->data->ty.as<PointerTypeNode>();
+ const auto* rhs_ptr = rhs->data->ty.as<PointerTypeNode>();
+ equal = lhs_ptr && rhs_ptr &&
+ ffi::StructuralEqual()(lhs_ptr->element_type,
rhs_ptr->element_type);
+ if (equal) {
+ equal_map_[lhs->data] = rhs->data;
+ }
+ }
+ equal = equal && lhs->dtype == rhs->dtype;
if (equal) {
rhs_buffer_map_[rhs] = lhs;
lhs_buffer_map_[lhs] = rhs;
diff --git a/src/target/source/codegen_c.cc b/src/target/source/codegen_c.cc
index 8bf390a508..f80a95469e 100644
--- a/src/target/source/codegen_c.cc
+++ b/src/target/source/codegen_c.cc
@@ -650,6 +650,12 @@ void CodeGenC::VisitExpr_(const NotNode* op, std::ostream&
os) { // NOLINT(*)
void CodeGenC::PrintCallExtern(Type ret_type, ffi::String global_symbol,
const ffi::Array<Expr>& args, bool
skip_first_arg,
std::ostream& os) { // NOLINT(*)
+ bool cast_pointer_return = ret_type.as<PointerTypeNode>();
+ if (cast_pointer_return) {
+ os << "((";
+ PrintType(ret_type, os);
+ os << ")";
+ }
os << global_symbol << "(";
for (size_t i = static_cast<size_t>(skip_first_arg); i < args.size(); ++i) {
this->PrintExpr(args[i], os);
@@ -658,6 +664,9 @@ void CodeGenC::PrintCallExtern(Type ret_type, ffi::String
global_symbol,
}
}
os << ")";
+ if (cast_pointer_return) {
+ os << ")";
+ }
}
void CodeGenC::VisitExpr_(const CallNode* op, std::ostream& os) { // NOLINT(*)
diff --git a/src/te/operation/create_primfunc.cc
b/src/te/operation/create_primfunc.cc
index 3d1b0fa0d3..d64b768e0e 100644
--- a/src/te/operation/create_primfunc.cc
+++ b/src/te/operation/create_primfunc.cc
@@ -641,10 +641,17 @@ Stmt GenerateStmtFromExternOp(const te::ExternOp&
extern_op, CreateFuncInfo* inf
for (int i = 0; i < extern_op->num_outputs(); ++i) {
const Buffer& placeholder = extern_op->output_placeholders[i];
const te::Tensor& output_tensor = extern_op.output(i);
- info->tensor2buffers[output_tensor] = placeholder;
+ Buffer output_buffer = placeholder;
if (!info->IsArg(output_tensor)) {
- info->root_alloc.push_back(placeholder);
+ PrimExpr zero_offset = IntImm(placeholder->elem_offset.ty(), 0);
+ if (const auto* offset_var = placeholder->elem_offset.as<VarNode>()) {
+ var_map[offset_var] = zero_offset;
+ }
+ output_buffer.CopyOnWrite()->elem_offset = zero_offset;
+ input_buffer_map[placeholder.get()] = output_buffer;
+ info->root_alloc.push_back(output_buffer);
}
+ info->tensor2buffers[output_tensor] = output_buffer;
}
// The access region does not need to be collected here, as it will
diff --git a/src/tirx/op/op.cc b/src/tirx/op/op.cc
index 363121933f..2344ba6137 100644
--- a/src/tirx/op/op.cc
+++ b/src/tirx/op/op.cc
@@ -562,6 +562,12 @@ PrimExpr reinterpret(PrimType t, PrimExpr value, Span
span) {
}
Expr reinterpret(Type target_ty, Expr value, Span span) {
+ if (value.as<StringImmNode>()) {
+ TVM_FFI_CHECK(target_ty.as<PointerTypeNode>(), TypeError)
+ << "String reinterpret requires a pointer target, but got " <<
target_ty;
+ return Call(std::move(target_ty), tirx::builtin::reinterpret(),
{std::move(value)}, {}, {},
+ std::move(span));
+ }
if (auto target_dtype = target_ty.as<PrimType>()) {
if (auto prim_value = value.as<PrimExpr>()) {
return reinterpret(target_dtype.value(), prim_value.value(),
std::move(span));
diff --git a/src/tirx/transform/lower_tvm_builtin.cc
b/src/tirx/transform/lower_tvm_builtin.cc
index 49ee7d306e..55790e51a4 100644
--- a/src/tirx/transform/lower_tvm_builtin.cc
+++ b/src/tirx/transform/lower_tvm_builtin.cc
@@ -553,6 +553,7 @@ class BuiltinLower : public StmtExprMutator {
int arg_type_index;
if (arg.as<StringImmNode>()) {
arg_type_index = ffi::TypeIndex::kTVMFFIRawStr;
+ arg = reinterpret(PointerType::VoidPointerTy(), std::move(arg));
} else if (arg->ty.as<PointerTypeNode>()) {
arg_type_index = IsArrayHandle(arg) ?
ffi::TypeIndex::kTVMFFIDLTensorPtr
: ffi::TypeIndex::kTVMFFIOpaquePtr;
diff --git a/tests/python/nightly/test_nnapi/__init__.py
b/tests/nightly/python/test_nnapi/__init__.py
similarity index 100%
rename from tests/python/nightly/test_nnapi/__init__.py
rename to tests/nightly/python/test_nnapi/__init__.py
diff --git a/tests/python/nightly/test_nnapi/conftest.py
b/tests/nightly/python/test_nnapi/conftest.py
similarity index 100%
rename from tests/python/nightly/test_nnapi/conftest.py
rename to tests/nightly/python/test_nnapi/conftest.py
diff --git a/tests/python/nightly/test_nnapi/infrastructure.py
b/tests/nightly/python/test_nnapi/infrastructure.py
similarity index 100%
rename from tests/python/nightly/test_nnapi/infrastructure.py
rename to tests/nightly/python/test_nnapi/infrastructure.py
diff --git a/tests/python/nightly/test_nnapi/test_from_exported_to_cuda.py
b/tests/nightly/python/test_nnapi/test_from_exported_to_cuda.py
similarity index 100%
rename from tests/python/nightly/test_nnapi/test_from_exported_to_cuda.py
rename to tests/nightly/python/test_nnapi/test_from_exported_to_cuda.py
diff --git a/tests/python/nightly/test_nnapi/test_network.py
b/tests/nightly/python/test_nnapi/test_network.py
similarity index 100%
rename from tests/python/nightly/test_nnapi/test_network.py
rename to tests/nightly/python/test_nnapi/test_network.py
diff --git a/tests/python/nightly/test_nnapi/test_ops.py
b/tests/nightly/python/test_nnapi/test_ops.py
similarity index 100%
rename from tests/python/nightly/test_nnapi/test_ops.py
rename to tests/nightly/python/test_nnapi/test_ops.py
diff --git a/tests/python/codegen/test_target_codegen_c_host.py
b/tests/python/codegen/test_target_codegen_c_host.py
index 5dac50d48e..8af6be0b79 100644
--- a/tests/python/codegen/test_target_codegen_c_host.py
+++ b/tests/python/codegen/test_target_codegen_c_host.py
@@ -227,5 +227,23 @@ def test_subroutine_call():
)
+def test_workspace_allocation_cast():
+ @I.ir_module
+ class Module:
+ @T.prim_func
+ def main(A: T.Buffer((256,), "float32")):
+ workspace = T.alloc_buffer((256,), "float32", scope="global")
+ for i in range(256):
+ workspace[i] = A[i]
+ for i in range(256):
+ A[i] = workspace[i]
+
+ built = tvm.tirx.build(Module, target="c")
+ assert "((float*)TVMBackendAllocWorkspace(" in built.inspect_source()
+
+ temp = utils.tempdir()
+ built.export_library(temp.relpath("workspace.so"))
+
+
if __name__ == "__main__":
tvm.testing.main()
diff --git a/tests/python/conftest.py b/tests/python/conftest.py
index 5d55ba5c6a..b3a7dc8378 100644
--- a/tests/python/conftest.py
+++ b/tests/python/conftest.py
@@ -14,17 +14,16 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-"""Configure pytest"""
+"""Configure pytest for TVM's Python test suite."""
-import sys
+import os
+from pathlib import Path
-COLLECT_IGNORE = []
-if sys.platform.startswith("win"):
- COLLECT_IGNORE.append("frontend/coreml")
- COLLECT_IGNORE.append("frontend/keras")
- COLLECT_IGNORE.append("frontend/pytorch")
- COLLECT_IGNORE.append("frontend/tensorflow")
- COLLECT_IGNORE.append("frontend/tflite")
- COLLECT_IGNORE.append("frontend/onnx")
- COLLECT_IGNORE.append("tir_base/test_tir_intrin.py")
+def pytest_sessionstart():
+ if os.getenv("CI", "") == "true":
+ from tvm.testing.utils import (
+ install_request_hook, # pylint: disable=import-outside-toplevel
+ )
+
+ install_request_hook(Path(__file__).with_name("request_hook.py"))
diff --git a/tests/python/disco/test_callback.py
b/tests/python/disco/test_callback.py
index 5d3d7f14b0..817e2289f6 100644
--- a/tests/python/disco/test_callback.py
+++ b/tests/python/disco/test_callback.py
@@ -25,13 +25,16 @@ import pytest
import tvm
import tvm.testing
+from tvm.script import ir as I
from tvm.script import relax as R
from tvm.script import tirx as T
from tvm.testing import env
@pytest.mark.gpu
[email protected](tvm.runtime.disco is None, reason="disco runtime is not
available")
@pytest.mark.skipif(not env.has_nccl(), reason="need nccl")
[email protected](not env.has_multi_gpu(), reason="need multiple gpus")
def test_callback():
"""Simulate lazy loading of parameters in a callback
@@ -39,22 +42,54 @@ def test_callback():
callback to load the parameters.
"""
- @R.function
- def transform_params(
- rank_arg: R.Prim(value="rank"),
- fget_item: R.Callable([R.Any, R.Prim("int64")], R.Any),
- ):
- rank = T.int64()
-
- A = fget_item(R.str("A"), R.prim_value(0))
- A = R.match_cast(A, R.Tensor([4, 4], "int32"))
- A = R.strided_slice(A, axes=[0], begin=[rank * 2], end=[(rank + 1) *
2])
-
- B = fget_item(R.str("B"), R.prim_value(1))
- B = R.match_cast(B, R.Tensor([2, 2], "float32"))
- B = R.strided_slice(B, axes=[1], begin=[rank * 1], end=[(rank + 1) *
1])
-
- return (A, B)
+ @I.ir_module(s_tir=True)
+ class Module:
+ @T.prim_func(private=True, s_tir=True)
+ def slice_A(
+ A: T.Buffer((4, 4), "int32"),
+ rank: T.int64,
+ A_sharded: T.Buffer((2, 4), "int32"),
+ ):
+ for i, j in T.grid(2, 4):
+ with T.sblock("slice_A"):
+ vi, vj = T.axis.remap("SS", [i, j])
+ A_sharded[vi, vj] = A[rank * 2 + vi, vj]
+
+ @T.prim_func(private=True, s_tir=True)
+ def slice_B(
+ B: T.Buffer((2, 2), "float32"),
+ rank: T.int64,
+ B_sharded: T.Buffer((2, 1), "float32"),
+ ):
+ for i in range(2):
+ with T.sblock("slice_B"):
+ vi = T.axis.spatial(2, i)
+ B_sharded[vi, 0] = B[vi, rank]
+
+ @R.function
+ def transform_params(
+ rank_arg: R.Prim("int64"),
+ fget_item: R.Callable([R.Any, R.Prim("int64")], R.Any),
+ ):
+ cls = Module
+
+ A = fget_item(R.str("A"), R.prim_value(0))
+ A = R.match_cast(A, R.Tensor([4, 4], "int32"))
+ A = R.call_tir(
+ cls.slice_A,
+ (A, rank_arg),
+ out_ty=R.Tensor([2, 4], "int32"),
+ )
+
+ B = fget_item(R.str("B"), R.prim_value(1))
+ B = R.match_cast(B, R.Tensor([2, 2], "float32"))
+ B = R.call_tir(
+ cls.slice_B,
+ (B, rank_arg),
+ out_ty=R.Tensor([2, 1], "float32"),
+ )
+
+ return (A, B)
pipeline = tvm.ir.transform.Sequential(
[
@@ -65,7 +100,7 @@ def test_callback():
)
with tvm.target.Target("cuda"):
- mod = tvm.IRModule.from_expr(transform_params)
+ mod = Module
mod = pipeline(mod)
built = tvm.compile(mod, "cuda")
diff --git a/tests/python/disco/test_ccl.py b/tests/python/disco/test_ccl.py
index 29d46d4ec1..82dbc27ab7 100644
--- a/tests/python/disco/test_ccl.py
+++ b/tests/python/disco/test_ccl.py
@@ -31,6 +31,9 @@ from tvm.runtime.vm import VirtualMachine
from tvm.s_tir import dlight as dl
from tvm.script import relax as R
+if di is None:
+ pytest.skip("disco runtime is not available", allow_module_level=True)
+
_all_session_kinds = [di.ThreadedSession, di.ProcessSession]
_compiled_ccl = get_global_func("runtime.disco.compiled_ccl",
allow_missing=True)
if _compiled_ccl is None:
diff --git a/tests/python/disco/test_custom_allreduce.py
b/tests/python/disco/test_custom_allreduce.py
index f49da02836..db23257703 100644
--- a/tests/python/disco/test_custom_allreduce.py
+++ b/tests/python/disco/test_custom_allreduce.py
@@ -26,7 +26,9 @@ from tvm_ffi import Shape
import tvm
import tvm.testing
from tvm.runtime import DataType, disco
-from tvm.runtime.disco import Session
+
+if disco is None:
+ pytest.skip("disco runtime is not available", allow_module_level=True)
class AllReduceStrategyType(enum.IntEnum):
@@ -56,7 +58,7 @@ _ccl = [ccl for ccl in _compiled_ccl() if ccl == "nccl"]
@pytest.mark.parametrize("strategy", _strategies)
def test_allreduce(shape, ccl, strategy):
devices = [0, 1]
- sess: Session = disco.ProcessSession(num_workers=len(devices))
+ sess = disco.ProcessSession(num_workers=len(devices))
sess.init_ccl(ccl, *devices)
num_elements = reduce(lambda x, y: x * y, shape)
diff --git a/tests/python/relax/test_analysis_type_analysis.py
b/tests/python/relax/test_analysis_type_analysis.py
index c98f5bd9ca..a367de212f 100644
--- a/tests/python/relax/test_analysis_type_analysis.py
+++ b/tests/python/relax/test_analysis_type_analysis.py
@@ -636,57 +636,13 @@ def _generate_prim_test_cases():
# LCA of a PrimType with itself yields itself
yield (R.Prim(dtype), R.Prim(dtype), R.Prim(dtype))
- # The LCA of two values, each statically known to be the same
- # value, is known to have that value.
- yield (
- R.Prim(value=tirx.const(0, dtype)),
- R.Prim(value=tirx.const(0, dtype)),
- R.Prim(value=tirx.const(0, dtype)),
- )
-
- # The LCA of two values, each of which is statically known to
- # have a different value, no longer knows the contained value.
- yield (
- R.Prim(value=tirx.const(0, dtype)),
- R.Prim(value=tirx.const(1, dtype)),
- R.Prim(dtype=dtype),
- )
-
- # LCA of a known variable with itself yields itself
- var_N = tirx.Var("N", dtype)
- yield (R.Prim(value=var_N), R.Prim(value=var_N), R.Prim(value=var_N))
-
- # LCA of a known variable with a known static value is no
- # longer known to have a specific value.
- yield (R.Prim(value=var_N), R.Prim(value=tirx.const(0, dtype)),
R.Prim(dtype=dtype))
- yield (R.Prim(value=tirx.const(0, dtype)), R.Prim(value=var_N),
R.Prim(dtype=dtype))
-
- var_M = tirx.Var("M", dtype)
- yield (R.Prim(value=var_N), R.Prim(value=var_M), R.Prim(dtype=dtype))
-
for dtype_a in dtypes:
for dtype_b in dtypes:
if dtype_a != dtype_b:
- # Unlike R.Tensor, R.Prim does not currently support a
- # value with an unknown datatype. If the dtype
- # differs between the two annotations, the next wider
- # category is R.Any.
+ # If the dtype differs between the two annotations,
+ # the next wider category is R.Any.
yield (R.Prim(dtype_a), R.Prim(dtype_b), R.Any)
- # Because the dtypes are different, even `R.Prim` containing
- # the same value in different representations (e.g.
- # `T.float32(0)` vs `T.float16(0)`) fall back to `R.Any`.
- yield (
- R.Prim(value=tirx.const(0, dtype_a)),
- R.Prim(value=tirx.const(0, dtype_b)),
- R.Any,
- )
-
- # And the same is true for known variable values
- var_N = tirx.Var("N", dtype_a)
- var_M = tirx.Var("M", dtype_b)
- yield (R.Prim(value=var_N), R.Prim(value=var_M), R.Any)
-
@pytest.mark.parametrize("test_case", list(_generate_prim_test_cases()))
def test_prim_type_lca(test_case):
@@ -806,7 +762,7 @@ def test_collect_nonnegative_expressions():
A: R.Tensor([1024, "M", "N-2"]),
B: R.Tensor([128, "N", "M+2"]),
C: R.Shape(["M", "N"]),
- D: R.Prim(value="N"),
+ D: R.Prim("int64"),
):
return R.tuple()
diff --git a/tests/python/relax/test_backend_transform_shape_lower.py
b/tests/python/relax/test_backend_transform_shape_lower.py
index d38e5dc8ab..746edf9ace 100644
--- a/tests/python/relax/test_backend_transform_shape_lower.py
+++ b/tests/python/relax/test_backend_transform_shape_lower.py
@@ -16,8 +16,6 @@
# under the License.
# ruff: noqa: F841
-import pytest
-
import tvm.script
import tvm.testing
from tvm import relax
@@ -817,84 +815,5 @@ def test_check_weights_with_dynamic_shape():
assert_structural_equal(after, expected)
[email protected](reason="value-bearing R.Prim annotations were removed")
-def test_update_symbolic_vars_in_match_cast_rhs():
- """Symbolic variables may be used on the RHS of match_cast"""
-
- @I.ir_module
- class Before:
- @R.function
- def main(
- arg_prim_value: R.Prim(value="n"),
- ):
- R.func_attr({"relax.force_pure": True})
- n = T.int64()
- shape = R.shape([n])
- m = T.int64()
- _ = R.match_cast(shape, R.Shape([m]))
- return R.prim_value(m)
-
- @I.ir_module
- class Expected:
- @R.function
- def main(arg_prim_value: R.Prim(value="n")) -> R.Prim("int64"):
- R.func_attr({"relax.force_pure": True})
- n = T.int64()
-
- shape_heap = R.call_builtin_with_ctx(
- "vm.builtin.alloc_shape_heap",
- [2],
- ty_args=(R.Tensor(dtype="int64", ndim=1),),
- )
- _ = R.call_packed(
- "vm.builtin.check_prim_value_info",
- arg_prim_value,
- R.dtype("int64"),
- "",
- ty_args=[R.Tuple],
- )
- _ = R.call_packed(
- "vm.builtin.match_prim_value",
- arg_prim_value,
- shape_heap,
- MatchShapeCode.STORE_TO_HEAP,
- 0,
- "",
- ty_args=[R.Tuple],
- )
- shape = R.call_packed(
- "vm.builtin.make_shape",
- shape_heap,
- 1,
- MakeShapeCode.LOAD_SHAPE,
- 0,
- ty_args=[R.Shape(ndim=1)],
- )
- _ = R.call_packed(
- "vm.builtin.match_shape",
- shape,
- shape_heap,
- 1,
- MatchShapeCode.STORE_TO_HEAP,
- 1,
- "",
- ty_args=[R.Tuple],
- )
-
- m = T.int64()
- _ = R.match_cast(shape, R.Shape([m]))
- gv = R.call_packed(
- "vm.builtin.make_prim_value",
- shape_heap,
- MakeShapeCode.LOAD_SHAPE,
- 1,
- ty_args=[R.Prim(value=m)],
- )
- return gv
-
- After = relax.transform.VMShapeLower(emit_err_ctx=False)(Before)
- assert_structural_equal(Expected, After)
-
-
if __name__ == "__main__":
tvm.testing.main()
diff --git a/tests/python/relax/test_bind_params.py
b/tests/python/relax/test_bind_params.py
index 14eddcb360..8d24e95d6c 100644
--- a/tests/python/relax/test_bind_params.py
+++ b/tests/python/relax/test_bind_params.py
@@ -112,26 +112,16 @@ prim_value_dtype = tvm.testing.parameter("int64",
"int32", "float32")
def test_bind_prim_value(prim_value_dtype):
- if prim_value_dtype != "int64":
- pytest.xfail(reason="Currently, only support int64 as known symbolic
value")
-
- N = tirx.Var("N", prim_value_dtype)
+ prim_type = tvm.ir.PrimType(prim_value_dtype)
+ param = relax.Var("A", prim_type)
+ before = relax.Function([param], param,
prim_type).with_attr("global_symbol", "main")
value = tirx.const(16, prim_value_dtype)
- @R.function
- def before(A: R.Prim(value=N)) -> R.Prim(value=N):
- R.func_attr({"global_symbol": "main"})
- B: R.Prim(value=N) = A
- return B
-
- @R.function
- def expected() -> R.Prim(value=value):
- R.func_attr({"global_symbol": "main"})
- return value
-
after = before.bind_params({"A": value})
- tvm.ir.assert_structural_equal(expected, after)
+ assert not after.params
+ tvm.ir.assert_structural_equal(after.ret_ty, prim_type)
+ tvm.ir.assert_structural_equal(after.body.body, value)
def test_error_on_unknown_var():
diff --git a/tests/python/relax/test_bind_symbolic_vars.py
b/tests/python/relax/test_bind_symbolic_vars.py
index bb494747b1..1f80917860 100644
--- a/tests/python/relax/test_bind_symbolic_vars.py
+++ b/tests/python/relax/test_bind_symbolic_vars.py
@@ -204,73 +204,6 @@ def test_bind_symbolic_vars_in_shape_expr():
tvm.ir.assert_structural_equal(expected, after)
[email protected](reason="value-bearing R.Prim annotations were removed")
-def test_bind_defining_of_symbolic_vars_in_prim_value():
- """R.Prim may define symbolic variables
-
- This case is a bit odd, because it always results in a
- fully-constrained parameter at the relax level. After binding in
- this test case, we have a function that accepts three parameters,
- and the third parameter must always be the number 16.
-
- However, this provides the most consistent behavior with other
- uses of `relax.Function.bind_symbolic_vars`, which restricts the
- allowed values for each parameter, but does not alter the number
- of parameters. This is in contrast to the `BindParams` pass,
- which provides a known value for relax parameters, removing them
- from the function signature.
-
- This convention also prevents surprise changes to the function
- signature, such as shown in
- `test_bind_symbolic_vars_with_expr_in_prim_value`.
- """
-
- @R.function(private=True)
- def before(A: R.Tensor(["M * N"]), x: R.Prim(value="M"), y:
R.Prim(value="N")):
- M = T.int64()
- N = T.int64()
- B = R.call_dps_packed("dummy_func", [A], out_ty=R.Tensor([2 * M * N]))
- return B
-
- @R.function(private=True)
- def expected(A: R.Tensor(["M * 16"]), x: R.Prim(value="M"), y:
R.Prim(value=16)):
- B = R.call_dps_packed("dummy_func", [A], out_ty=R.Tensor([M * 32]))
- return B
-
- after = before.bind_symbolic_vars({"N": 16})
- tvm.ir.assert_structural_equal(expected, after)
-
-
-def test_bind_usage_of_symbolic_vars_in_prim_value():
- """R.Prim may use symbolic variables defined by other parameters
-
- Like test_bind_defining_of_symbolic_vars_in_prim_value, but with
- R.Prim using a symbolic variable rather than defining it.
-
- This also demonstrates why we should not remove fully-constrained
- R.Prim function parameters. In this case, we have a function that
- accepts two parameters, and we have specialized the shape of the
- first parameter. It would be unexpected for specialization of the
- first parameter to result in removal of a different parameter
- altogether.
- """
-
- @R.function(private=True)
- def before(A: R.Tensor(["M", "N"]), x: R.Prim(value="M*N")):
- M = T.int64()
- N = T.int64()
- B = R.call_dps_packed("dummy_func", [A], out_ty=R.Tensor([2 * M * N]))
- return B
-
- @R.function(private=True)
- def expected(A: R.Tensor([16, 16]), x: R.Prim(value=256)):
- B = R.call_dps_packed("dummy_func", [A], out_ty=R.Tensor([512]))
- return B
-
- after = before.bind_symbolic_vars({"M": 16, "N": 16})
- tvm.ir.assert_structural_equal(expected, after)
-
-
def test_bind_strided_slice():
"""relax.op.strided_slice stores Expr attributes"""
diff --git a/tests/python/relax/test_blockbuilder_core.py
b/tests/python/relax/test_blockbuilder_core.py
index 725ebd8219..dceadc3af5 100644
--- a/tests/python/relax/test_blockbuilder_core.py
+++ b/tests/python/relax/test_blockbuilder_core.py
@@ -662,8 +662,8 @@ def test_emit_nested_tuple(emit_nested_tuple):
@R.function
def func(
- n_1: R.Prim(value="n"),
- m_1: R.Prim(value="m"),
+ n_1: R.Prim("int64"),
+ m_1: R.Prim("int64"),
x: R.Tensor(("n", "m"), dtype="float32"),
y: R.Tensor(("m", "n"), dtype="float32"),
):
@@ -673,8 +673,8 @@ def test_emit_nested_tuple(emit_nested_tuple):
@R.function
def func(
- n_1: R.Prim(value="n"),
- m_1: R.Prim(value="m"),
+ n_1: R.Prim("int64"),
+ m_1: R.Prim("int64"),
x: R.Tensor(("n", "m"), dtype="float32"),
y: R.Tensor(("m", "n"), dtype="float32"),
):
diff --git a/tests/python/relax/test_dataflow_rewriter.py
b/tests/python/relax/test_dataflow_rewriter.py
index c7a9ae2cbb..8446f770ac 100644
--- a/tests/python/relax/test_dataflow_rewriter.py
+++ b/tests/python/relax/test_dataflow_rewriter.py
@@ -14,23 +14,16 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-# ruff: noqa: F841
import pytest
import tvm.testing
-from tvm import relax
-from tvm.ir import Op
from tvm.script import ir as I
from tvm.script import relax as R
from tvm.script import tirx as T
-def inspect_tensor_field(op_name, *args):
- return relax.Call(Op.get(f"relax.inspect.{op_name}"), args)
-
-
def test_rewrite_defined_by_ir_module():
@R.rewriter
class Rewriter:
@@ -372,148 +365,6 @@ def test_recursive_rewrite_rules():
tvm.ir.assert_structural_equal(expected, after)
[email protected](reason="value-bearing R.Prim match-cast semantics were
removed")
-def test_rewrite_of_arbitrary_dtype():
- """A pattern-match may apply to a tensor with unknown dtype
-
- In this test case, a pattern identifies `R.strided_slice` usage
- which returns the last slice of an array, and replaces it with a
- view into the input array.
-
- """
-
- @R.rewriter
- class Rewriter:
- @R.function
- def pattern(A: R.Tensor(["M", "N"])) -> R.Tensor(["N"]):
- M = T.int64()
- N = T.int64()
- last_slice_2d: R.Tensor([1, N]) = R.strided_slice(A, axes=[0],
begin=[M - 1], end=[M])
- last_slice_1d: R.Tensor([N]) = R.squeeze(last_slice_2d, axis=0)
- return last_slice_1d
-
- @R.function
- def replacement(A: R.Tensor(["M", "N"])) -> R.Tensor(["N"]):
- M = T.int64()
- N = T.int64()
-
- # A primitive-valued Relax inspect Call can be used in a
- # Relax context that accepts an Expr. Currently,
- # this requires `R.match_cast` to produce a TIR symbolic
- # variable from the Relax Expr.
- bits_per_element = T.uint8()
- _ = R.match_cast(
- inspect_tensor_field("tensor_dtype_bits", A),
- R.Prim(value=bits_per_element),
- )
- lanes_per_element = T.uint16()
- _ = R.match_cast(
- inspect_tensor_field("tensor_dtype_lanes", A),
- R.Prim(value=lanes_per_element),
- )
-
- last_slice = R.memory.view(
- A,
- [N],
- relative_byte_offset=(M - 1)
- * N
- * T.ceildiv(
- bits_per_element.astype("int64") *
lanes_per_element.astype("int64"), 8
- ),
- )
- return last_slice
-
- @I.ir_module
- class Before:
- @R.function
- def main(
- A: R.Tensor([32, 16], "float16"),
- B: R.Tensor(["P", "Q"], "int4x8"),
- C: R.Tensor([16, 32]),
- ):
- P = T.int64()
- Q = T.int64()
-
- A_slice_2d = R.strided_slice(A, axes=[0], begin=[31], end=[32])
- A_slice_1d = R.squeeze(A_slice_2d, axis=0)
-
- B_slice_2d = R.strided_slice(B, axes=[0], begin=[P - 1], end=[P])
- B_slice_1d = R.squeeze(B_slice_2d, axis=0)
-
- C_slice_2d = R.strided_slice(C, axes=[0], begin=[15], end=[16])
- C_slice_1d = R.squeeze(C_slice_2d, axis=0)
-
- return (A_slice_1d, B_slice_1d, C_slice_1d)
-
- @I.ir_module
- class Expected:
- @R.function
- def main(
- A: R.Tensor([32, 16], "float16"),
- B: R.Tensor(["P", "Q"], "int4x8"),
- C: R.Tensor([16, 32]),
- ):
- P = T.int64()
- Q = T.int64()
-
- # The pattern matches any 2-d tensor, with any data type.
- # When the match's shape and dtype are both known,
- # normalization and canonicalization produces a statically
- # known value for `relative_byte_offset`.
- #
- # Relative offset is `(31 rows) *
- # (16 elements/row) *
- # (2 bytes/element)`
- A_slice_1d = R.memory.view(A, shape=[16], relative_byte_offset=992)
-
- # The pattern can also match a 2-d tensor with dynamic
- # shape. The `relative_byte_offset` uses the known
- # datatype (4 bytes for each int4x8), but with dynamic
- # shape variables substituted in where required.
- #
- # Relative offset is `((P-1) rows) *
- # (Q elements/row) *
- # (4 bytes/element)`
- B_slice_1d = R.memory.view(B, shape=[Q], relative_byte_offset=(P -
1) * Q * 4)
-
- # The pattern can also match a 2-d tensor with static
- # shape, but unknown data type. The
- # `relative_byte_offset` is determined based on the known
- # number of elements, and the dynamic size of each
- # element.
- #
- # Relative offset is `(15 rows) *
- # (32 elements/row) *
- # (ceildiv(bits*lanes,8) bytes/element)`
- C_bits_per_element = T.uint8()
- C_bits_prim_value = inspect_tensor_field("tensor_dtype_bits", C)
- _ = R.match_cast(
- C_bits_prim_value,
- R.Prim(value=C_bits_per_element),
- )
- C_lanes_per_element = T.uint16()
- C_lanes_prim_value = inspect_tensor_field("tensor_dtype_lanes", C)
- _ = R.match_cast(
- C_lanes_prim_value,
- R.Prim(value=C_lanes_per_element),
- )
-
- C_slice_1d = R.memory.view(
- C,
- shape=[32],
- relative_byte_offset=(
- (C_bits_per_element.astype("int64") *
C_lanes_per_element.astype("int64") + 7)
- // 8
- )
- * 480,
- )
-
- return (A_slice_1d, B_slice_1d, C_slice_1d)
-
- after = Rewriter(Before)
- tvm.ir.assert_structural_equal(Expected, after)
-
-
def test_rewrite_may_introduce_private_relax_subroutines():
"""The replacement may contain subroutines"""
diff --git a/tests/python/relax/test_op_binary.py
b/tests/python/relax/test_op_binary.py
index f5d12bbe67..e3befa2a97 100644
--- a/tests/python/relax/test_op_binary.py
+++ b/tests/python/relax/test_op_binary.py
@@ -144,23 +144,6 @@ def
test_infer_ty_binary_arith_prim_value_with_prim_value(binary_arith_op: Calla
_check_inference(bb, binary_arith_op(x, y), tvm.ir.PrimType("float32"))
[email protected]("binary_arith_op,tir_arith_op", binary_arith_ops)
[email protected](reason="Not yet implemented")
-def test_infer_ty_binary_arith_known_prim_value_with_prim_value(
- binary_arith_op: Callable, tir_arith_op
-):
- bb = relax.BlockBuilder()
-
- tir_x = tirx.Var("tir_x", "float32")
- tir_y = tirx.Var("tir_y", "float32")
-
- x = relax.Var("x", R.Prim(value=tir_x))
- y = relax.Var("y", R.Prim(value=tir_y))
-
- _check_inference(bb, binary_arith_op(x, y), tvm.ir.PrimType("float32"))
- _check_inference(bb, binary_arith_op(y, x), tvm.ir.PrimType("float32"))
-
-
binary_cmp_ops = [
(relax.op.equal, tirx.EQ),
(relax.op.greater, tirx.GT),
@@ -206,21 +189,6 @@ def
test_infer_ty_binary_cmp_prim_value_to_prim_value(binary_cmp_op: Callable):
_check_inference(bb, binary_cmp_op(y, x), tvm.ir.PrimType("bool"))
[email protected]("binary_cmp_op,tir_cmp_op", binary_cmp_ops)
[email protected](reason="Not yet implemented")
-def test_infer_ty_binary_cmp_known_prim_value_to_prim_value(binary_cmp_op:
Callable, tir_cmp_op):
- bb = relax.BlockBuilder()
-
- tir_x = tirx.Var("tir_x", "float32")
- tir_y = tirx.Var("tir_y", "float32")
-
- x = relax.Var("x", R.Prim(value=tir_x))
- y = relax.Var("y", R.Prim(value=tir_y))
-
- _check_inference(bb, binary_cmp_op(x, y), tvm.ir.PrimType("bool"))
- _check_inference(bb, binary_cmp_op(y, x), tvm.ir.PrimType("bool"))
-
-
@pytest.mark.parametrize("binary_arith_op", [row[0] for row in
binary_arith_ops])
def test_binary_infer_ty_shape_symbolic(binary_arith_op: Callable):
bb = relax.BlockBuilder()
diff --git a/tests/python/relax/test_transform_compute_prim_value.py
b/tests/python/relax/test_transform_compute_prim_value.py
index a7b89f0654..37a4fafcb1 100644
--- a/tests/python/relax/test_transform_compute_prim_value.py
+++ b/tests/python/relax/test_transform_compute_prim_value.py
@@ -15,8 +15,6 @@
# specific language governing permissions and limitations
# under the License.
-import pytest
-
import tvm
import tvm.testing
from tvm.script import ir as I
@@ -84,34 +82,5 @@ def test_prim_value_in_branch_condition():
tvm.ir.assert_structural_equal(After, Expected)
[email protected](reason="value-bearing R.Prim annotations were removed")
-def test_prim_value_in_pure_function():
- @I.ir_module
- class Before:
- @R.function
- def main(_N: R.Prim(value="N"), _M: R.Prim(value="M")) ->
R.Prim(value="N*M"):
- N = T.int64()
- M = T.int64()
- out = R.prim_value(N * M)
- return out
-
- @I.ir_module
- class Expected:
- @R.function
- def main(_N: R.Prim(value="N"), _M: R.Prim(value="M")) ->
R.Prim(value="N*M"):
- N = T.int64()
- M = T.int64()
- out = Expected.compute_symbolic_expr(R.prim_value(N),
R.prim_value(M))
- return out
-
- @T.prim_func(private=True, s_tir=True)
- def compute_symbolic_expr(N: T.int64, M: T.int64) -> T.int64:
- T.func_attr({"tirx.is_host_func": True})
- T.ret(N * M)
-
- After = tvm.relax.transform.ComputePrimValue()(Before)
- tvm.ir.assert_structural_equal(After, Expected)
-
-
if __name__ == "__main__":
tvm.testing.main()
diff --git a/tests/python/relax/test_transform_lazy_transform_params.py
b/tests/python/relax/test_transform_lazy_transform_params.py
index 3d6a3ea88e..23b5d61c70 100644
--- a/tests/python/relax/test_transform_lazy_transform_params.py
+++ b/tests/python/relax/test_transform_lazy_transform_params.py
@@ -16,7 +16,6 @@
# under the License.
# ruff: noqa: F841
import numpy as np
-import pytest
import tvm
import tvm.testing
@@ -744,53 +743,6 @@ def test_params_without_tuple():
tvm.ir.assert_structural_equal(After, Expected)
[email protected](reason="value-bearing R.Prim annotations were removed")
-def test_retain_before_num_input():
- """Only lazily load parameters after num_input"""
-
- @I.ir_module(s_tir=True)
- class Before:
- @R.function
- def transform_params(
- relax_rank: R.Prim(value="rank"),
- A: R.Tensor([16, 16], "float32"),
- B: R.Tensor([16, 16], "float32"),
- ):
- R.func_attr({"num_input": 1})
- rank = T.int64()
- A_sharded = R.strided_slice(
- A, axes=[0], begin=[rank * 8], end=[(rank + 1) * 8],
assume_inbound=True
- )
- B_sharded = R.strided_slice(
- B, axes=[1], begin=[rank * 8], end=[(rank + 1) * 8],
assume_inbound=True
- )
- return (A_sharded, B_sharded)
-
- @I.ir_module(s_tir=True)
- class Expected:
- @R.function(pure=False)
- def transform_params(relax_rank: R.Prim(value="rank")):
- R.func_attr({"num_input": 1})
- rank = T.int64()
-
- A = R.call_packed("get_item", R.prim_value(0), ty_args=[R.Any])
- A = R.match_cast(A, R.Tensor([16, 16], "float32"))
- A_sharded = R.strided_slice(
- A, axes=[0], begin=[rank * 8], end=[(rank + 1) * 8],
assume_inbound=True
- )
-
- B = R.call_packed("get_item", R.prim_value(1), ty_args=[R.Any])
- B = R.match_cast(B, R.Tensor([16, 16], "float32"))
- B_sharded = R.strided_slice(
- B, axes=[1], begin=[rank * 8], end=[(rank + 1) * 8],
assume_inbound=True
- )
-
- return (A_sharded, B_sharded)
-
- After = LazyTransformParams(fset_item=None)(Before)
- tvm.ir.assert_structural_equal(After, Expected)
-
-
def test_params_without_tuple_with_symbolic_var():
@I.ir_module(s_tir=True)
class Before:
@@ -838,103 +790,6 @@ def test_get_item_callback():
tvm.ir.assert_structural_equal(After, Expected)
[email protected](reason="value-bearing R.Prim annotations were removed")
-def test_get_item_callback_num_attrs():
- @I.ir_module(s_tir=True)
- class Before:
- @R.function(pure=False)
- def transform_params(
- rank_arg: R.Prim(value="rank"),
- world_size_arg: R.Prim(value="world_size"),
- weight_A: R.Tensor([16, 64], "float32"),
- weight_B: R.Tensor([1024, 2048], "float32"),
- ):
- R.func_attr({"num_input": 2})
-
- rank = T.int64()
- world_size = T.int64()
-
- _ = R.assert_op(
- R.prim_value(16 % world_size == 0),
- [R.prim_value(16), R.prim_value(world_size)],
- format=(
- "World size must evenly divide A.shape[0] ({}), but
received world size of {}."
- ),
- )
- weight_A = R.strided_slice(
- weight_A,
- axes=[0],
- begin=[rank * 16 // world_size],
- end=[(rank + 1) * 16 // world_size],
- )
-
- _ = R.assert_op(
- R.prim_value(2048 % world_size == 0),
- [R.prim_value(2048), R.prim_value(world_size)],
- format=(
- "World size must evenly divide B.shape[1] ({}), but
received world size of {}."
- ),
- )
- weight_B = R.strided_slice(
- weight_B,
- axes=[1],
- begin=[rank * 2048 // world_size],
- end=[(rank + 1) * 2048 // world_size],
- )
-
- return (weight_A, weight_B)
-
- @I.ir_module(s_tir=True)
- class Expected:
- @R.function(pure=False)
- def transform_params(
- rank_arg: R.Prim(value="rank"),
- world_size_arg: R.Prim(value="world_size"),
- fget_item: R.Callable([R.Prim("int64"), R.Any], R.Any),
- ):
- R.func_attr({"num_input": 3})
-
- rank = T.int64()
- world_size = T.int64()
-
- _ = R.assert_op(
- R.prim_value(16 % world_size == 0),
- [R.prim_value(16), R.prim_value(world_size)],
- format=(
- "World size must evenly divide A.shape[0] ({}), but
received world size of {}."
- ),
- )
- weight_A = fget_item(R.prim_value(0), R.str("weight_A"))
- weight_A = R.match_cast(weight_A, R.Tensor([16, 64], "float32"))
- weight_A = R.strided_slice(
- weight_A,
- axes=[0],
- begin=[rank * 16 // world_size],
- end=[(rank + 1) * 16 // world_size],
- )
-
- _ = R.assert_op(
- R.prim_value(2048 % world_size == 0),
- [R.prim_value(2048), R.prim_value(world_size)],
- format=(
- "World size must evenly divide B.shape[1] ({}), but
received world size of {}."
- ),
- )
- weight_B = fget_item(R.prim_value(1), R.str("weight_B"))
- weight_B = R.match_cast(weight_B, R.Tensor([1024, 2048],
"float32"))
- weight_B = R.strided_slice(
- weight_B,
- axes=[1],
- begin=[rank * 2048 // world_size],
- end=[(rank + 1) * 2048 // world_size],
- )
-
- return (weight_A, weight_B)
-
- After = relax.transform.LazyGetInput()(Before)
- tvm.ir.assert_structural_equal(After, Expected)
-
-
def test_get_item_callback_dynamic_shape():
@I.ir_module(s_tir=True)
class Before:
diff --git a/tests/python/relax/test_transform_lift_transform_params.py
b/tests/python/relax/test_transform_lift_transform_params.py
index 9bc032b831..9cb73c79e1 100644
--- a/tests/python/relax/test_transform_lift_transform_params.py
+++ b/tests/python/relax/test_transform_lift_transform_params.py
@@ -1684,8 +1684,8 @@ def
test_symbolic_var_defined_in_params_but_used_in_weights():
In order to be a source of definition, a symbolic variable in the
parameters must occur as a distinct parameter, as a tensor shape
- `R.Tensor(["var"])`, an explicit `R.Shape(["var"])`, or as a
- `R.Prim(value="var")`. A variable that is part of a larger
+ `R.Tensor(["var"])` or an explicit `R.Shape(["var"])`. A variable
+ that is part of a larger
expression, such as `R.Tensor(["m * n"])`, are variable usages,
not variable definitions.
"""
diff --git a/tests/python/relax/test_transform_remove_unused_parameters.py
b/tests/python/relax/test_transform_remove_unused_parameters.py
index 3eaf8270bd..7e500c5204 100644
--- a/tests/python/relax/test_transform_remove_unused_parameters.py
+++ b/tests/python/relax/test_transform_remove_unused_parameters.py
@@ -63,11 +63,9 @@ def test_replace_symbolic_variables():
removing the `R.Tensor` argument, we may need to provide
additional parameters to define the symbolic variables.
- Value-bearing `R.Prim(value=...)` annotations were removed in the tirx
- refactor (a `PrimType` carries only a dtype and defines no TIR var, which
- leaves the var undefined under the strict tirx well-formedness verifier).
- The replacement is to promote each free symbolic variable through a 1-D
- `R.Shape` parameter, which actually *defines* the variable.
+ A `PrimType` carries only a dtype and defines no TIR var. Each free
+ symbolic variable is therefore promoted through a 1-D `R.Shape`
+ parameter, which actually *defines* the variable.
"""
@I.ir_module
@@ -131,12 +129,11 @@ def test_no_extra_symbolic_variables():
tvm.ir.assert_structural_equal(After, Expected)
-def test_remove_extra_prim_variables():
- """Remove parameters that only serve to define existing symbolic variables
+def test_remove_extra_prim_parameters():
+ """Remove unused scalar parameters.
- If a `R.Prim` parameter provies a definition of a symbolic
- variable, but that symbolic variable can be determined from a
- different parameter, then the `R.Prim` parameter can be removed.
+ The tensor parameter already defines the symbolic dimensions, while the
+ dtype-only scalar parameters are unused by the private function.
"""
@I.ir_module
@@ -149,7 +146,9 @@ def test_remove_extra_prim_variables():
@R.function(private=True)
def func(
- A: R.Tensor(["m", "n"], "float32"), _m: R.Prim(value="m"), _n:
R.Prim(value="n")
+ A: R.Tensor(["m", "n"], "float32"),
+ _m: R.Prim("int64"),
+ _n: R.Prim("int64"),
) -> R.Tensor(["m", "n"], "float32"):
m = T.int64()
n = T.int64()
diff --git a/tests/python/relax/test_tvmscript_parser.py
b/tests/python/relax/test_tvmscript_parser.py
index 7f7311ba60..1ca64a44ad 100644
--- a/tests/python/relax/test_tvmscript_parser.py
+++ b/tests/python/relax/test_tvmscript_parser.py
@@ -1313,6 +1313,14 @@ def test_scalar_tensor_as_branch_condition():
tvm.ir.assert_structural_equal(if_else.cond.ty, R.Tensor([], "bool"))
+def test_prim_annotation_requires_dtype():
+ with pytest.raises(TypeError, match="missing 1 required positional
argument: 'dtype'"):
+ R.Prim()
+
+ with pytest.raises(TypeError, match="unexpected keyword argument 'value'"):
+ R.Prim(value="n")
+
+
def test_prim_value_as_branch_condition():
"""In addition to scalar tensor, can use R.Prim condition"""
diff --git a/tests/python/relax/test_utils.py b/tests/python/relax/test_utils.py
index b654edfc31..d87b1a761b 100644
--- a/tests/python/relax/test_utils.py
+++ b/tests/python/relax/test_utils.py
@@ -24,7 +24,6 @@ import tvm
from tvm import relax
from tvm.ir.base import assert_structural_equal
from tvm.script.parser import relax as R
-from tvm.script.parser import tirx as T
def test_copy_with_new_vars():
@@ -171,7 +170,6 @@ def test_structural_equal_of_call_nodes():
tvm.ir.assert_structural_equal(uses_same_object_twice,
uses_two_different_objects)
[email protected](reason="value-bearing R.Prim annotations were removed")
def test_structural_equal_with_recursive_lambda_function():
"""A recursive lambda function may be checked for structural equality
@@ -188,15 +186,14 @@ def
test_structural_equal_with_recursive_lambda_function():
@R.function
def func(n: R.Prim("int64")):
@R.function
- def recursive_lambda(i_arg: R.Prim(value="i")) -> R.Prim("int64"):
- i = T.int64()
- if R.prim_value(i == 0):
- output = R.prim_value(T.int64(0))
+ def recursive_lambda(i_arg: R.Prim("int64")) -> R.Prim("int64"):
+ condition = R.equal(i_arg, R.prim_value(0))
+ if condition:
+ output = R.prim_value(0)
else:
- remainder_relax = recursive_lambda(R.prim_value(i - 1))
- remainder_tir = T.int64()
- _ = R.match_cast(remainder_relax,
R.Prim(value=remainder_tir))
- output = R.prim_value(i + remainder_tir)
+ next_i = R.subtract(i_arg, R.prim_value(1))
+ remainder = recursive_lambda(next_i)
+ output = R.add(i_arg, remainder)
return output
return recursive_lambda(n)
@@ -219,17 +216,16 @@ def
test_structural_equal_with_distinct_recursive_lambda_function():
@R.function(private=True)
def func_a(n: R.Prim("int64")):
@R.function
- def recursive_lambda(i_arg: R.Prim(value="i")) -> R.Prim("int64"):
- i = T.int64()
- if R.prim_value(i == 0):
- output = R.prim_value(T.int64(0))
- # ^
- # The first mismatch is here ^
+ def recursive_lambda(i_arg: R.Prim("int64")) -> R.Prim("int64"):
+ condition = R.equal(i_arg, R.prim_value(0))
+ if condition:
+ output = R.prim_value(0)
+ # ^
+ # The first mismatch is here
else:
- remainder_relax = recursive_lambda(R.prim_value(i - 1))
- remainder_tir = T.int64()
- _ = R.match_cast(remainder_relax, R.Prim(value=remainder_tir))
- output = R.prim_value(i + remainder_tir)
+ next_i = R.subtract(i_arg, R.prim_value(1))
+ remainder = recursive_lambda(next_i)
+ output = R.add(i_arg, remainder)
return output
return recursive_lambda(n)
@@ -237,17 +233,16 @@ def
test_structural_equal_with_distinct_recursive_lambda_function():
@R.function(private=True)
def func_b(n: R.Prim("int64")):
@R.function
- def recursive_lambda(i_arg: R.Prim(value="i")) -> R.Prim("int64"):
- i = T.int64()
- if R.prim_value(i == 0):
- output = R.prim_value(T.int64(1))
- # ^
- # The first mismatch is here ^
+ def recursive_lambda(i_arg: R.Prim("int64")) -> R.Prim("int64"):
+ condition = R.equal(i_arg, R.prim_value(0))
+ if condition:
+ output = R.prim_value(1)
+ # ^
+ # The first mismatch is here
else:
- remainder_relax = recursive_lambda(R.prim_value(i - 1))
- remainder_tir = T.int64()
- _ = R.match_cast(remainder_relax, R.Prim(value=remainder_tir))
- output = R.prim_value(i * remainder_tir)
+ next_i = R.subtract(i_arg, R.prim_value(1))
+ remainder = recursive_lambda(next_i)
+ output = R.multiply(i_arg, remainder)
return output
return recursive_lambda(n)
@@ -262,10 +257,11 @@ def
test_structural_equal_with_distinct_recursive_lambda_function():
"value",
"body",
"blocks[0]",
- "bindings[0]",
+ "bindings[1]",
+ "value",
+ "true_branch",
+ "body",
"value",
- "cond",
- "a",
]
with pytest.raises(ValueError, match=re.escape(".".join(mismatch_path))):
diff --git a/tests/python/relax/test_vm_build.py
b/tests/python/relax/test_vm_build.py
index 304ae33d77..5156202726 100644
--- a/tests/python/relax/test_vm_build.py
+++ b/tests/python/relax/test_vm_build.py
@@ -559,100 +559,6 @@ def test_vm_relax_symbolic_shape_tuple(exec_mode):
func(R.prim_value(2))
[email protected](reason="value-bearing R.Prim annotations are erased to
dtype-only PrimType")
-def test_vm_relax_symbolic_prim_value(exec_mode):
- @I.ir_module(s_tir=True)
- class mod:
- @R.function
- def main(shape: R.Prim(value="n")):
- n = T.int64()
- return R.prim_value(n * n)
-
- target = tvm.target.Target("llvm", host="llvm")
- ex = relax.build(mod, target, exec_mode=exec_mode)
- vm = relax.VirtualMachine(ex, tvm.cpu())
-
- func = vm["main"]
-
- assert func(2) == 4
-
- with pytest.raises(TypeError):
- func(Shape([2]))
-
-
[email protected](reason="value-bearing R.Prim annotations are erased to
dtype-only PrimType")
-def test_vm_relax_multiple_symbolic_prim_value(exec_mode):
- """Like test_vm_relax_symbolic_prim_value, but with multiple variables"""
-
- @I.ir_module(s_tir=True)
- class mod:
- @R.function
- def main(
- # Provides definition of "n"
- _n: R.Prim(value="n"),
- # Requires definitions of both "n" and "m", but cannot
- # provide either.
- _shape: R.Shape(["n*2", "m*2"]),
- # Provides definition of "m"
- _m: R.Prim(value="m"),
- ):
- n = T.int64()
- m = T.int64()
- return R.shape([n * n, m + 1])
-
- target = tvm.target.Target("llvm", host="llvm")
- ex = relax.build(mod, target, exec_mode=exec_mode)
- vm = relax.VirtualMachine(ex, tvm.cpu())
-
- func = vm["main"]
-
- assert func(2, Shape([4, 12]), 6) == (4, 7)
-
- with pytest.raises(RuntimeError):
- func(2, Shape([4, 12]), 1)
-
- with pytest.raises(RuntimeError):
- func(Shape([2]))
-
-
[email protected](reason="Current support for R.Prim with known value is
primarily for int64")
[email protected]("exec_mode", EXEC_MODE)
-def test_vm_relax_prim_value_fp32(exec_mode):
- """A Expr may be R.prim('float32')
-
- Unlike shape tuples, which must contain int64, a Expr may be
- any type that can be represented as a single primitive value.
- """
-
- @I.ir_module(s_tir=True)
- class mod:
- @R.function
- def main(
- # First failure occurs during parsing. The syntactic
- # sugar for symbolic variables assumes that all symbolic
- # variables are int64, rather than using the type that is
- # later declared.
- _x: R.Prim(value="half_fill_value"),
- ):
- half_fill_value = T.float32()
- # Second failure occurs when calling `relax.op.full`. The
- # `fill_value` is expected to be a scalar constant
- # (R.Tensor with 0-dim shape), not a primitive value, even
- # though these are semantically the same.
- return R.full(shape=[16, 16], fill_value=R.prim_value(2 *
half_fill_value))
-
- target = tvm.target.Target("llvm", host="llvm")
- # Third failure occurs here. The current codegen assumes that all
- # symbolic variables are int64.
- ex = relax.build(mod, target, exec_mode=exec_mode)
- vm = relax.VirtualMachine(ex, tvm.cpu())
-
- func = vm["main"]
-
- res = func(16.0).numpy()
- assert np.all(res == 32.0)
-
-
def test_vm_relax_dyn_tir_shape(exec_mode):
# case where TIR variables are unbound in generated PrimFunc
bb = relax.BlockBuilder()
diff --git a/tests/scripts/request_hook/request_hook.py
b/tests/python/request_hook.py
similarity index 92%
rename from tests/scripts/request_hook/request_hook.py
rename to tests/python/request_hook.py
index 48ac5e2a30..f3faf77a19 100644
--- a/tests/scripts/request_hook/request_hook.py
+++ b/tests/python/request_hook.py
@@ -35,7 +35,7 @@ class TvmRequestHook(urllib.request.Request):
msg = (
f"Uncaught URL found in CI: {url}. "
"Avoid network access or arrange a stable project-managed
mirror, "
- "then add it to URL_MAP in
tests/scripts/request_hook/request_hook.py."
+ "then add it to URL_MAP in tests/python/request_hook.py."
)
raise RuntimeError(msg)
@@ -46,6 +46,8 @@ class TvmRequestHook(urllib.request.Request):
def init():
global LOGGER
+ if urllib.request.Request is TvmRequestHook:
+ return
urllib.request.Request = TvmRequestHook
LOGGER = logging.getLogger("tvm_request_hook")
LOGGER.setLevel(logging.DEBUG)
diff --git
a/tests/python/s_tir/transform/test_s_tir_transform_plan_update_buffer_allocation_location.py
b/tests/python/s_tir/transform/test_s_tir_transform_plan_update_buffer_allocation_location.py
index d5173bcc13..9c59c321be 100644
---
a/tests/python/s_tir/transform/test_s_tir_transform_plan_update_buffer_allocation_location.py
+++
b/tests/python/s_tir/transform/test_s_tir_transform_plan_update_buffer_allocation_location.py
@@ -384,9 +384,13 @@ def test_dltensor_buffer_is_unlowered():
@T.prim_func(s_tir=True)
def before(dlpack_handle: T.handle, axis: T.int64) -> T.int64:
ndim: T.int32 = T.tvm_struct_get(dlpack_handle, 0, 5, "int32")
- stride_ptr: T.let[T.handle("int64")] = T.tvm_struct_get(dlpack_handle,
0, 4, "handle")
+ stride_ptr: T.let[T.handle("int64")] = T.tvm_struct_get(
+ dlpack_handle, 0, 4, dtype=T.handle("int64").ty
+ )
if T.isnullptr(stride_ptr):
- shape_ptr: T.let[T.handle("int64")] =
T.tvm_struct_get(dlpack_handle, 0, 3, "handle")
+ shape_ptr: T.let[T.handle("int64")] = T.tvm_struct_get(
+ dlpack_handle, 0, 3, dtype=T.handle("int64").ty
+ )
shape = T.decl_buffer(ndim, "int64", data=shape_ptr)
product = T.decl_buffer([], "int64")
product[()] = 1
diff --git a/tests/python/te/test_te_create_primfunc.py
b/tests/python/te/test_te_create_primfunc.py
index 90d7b69a55..48f9c1bf7b 100644
--- a/tests/python/te/test_te_create_primfunc.py
+++ b/tests/python/te/test_te_create_primfunc.py
@@ -293,6 +293,43 @@ def test_extern():
_check_workload(te_extern, tir_extern)
+def te_extern_epilogue():
+ A = te.placeholder((4, 3), name="A")
+ B = te.placeholder((3, 2), name="B")
+ C = te.extern(
+ (4, 2),
+ [A, B],
+ lambda ins, outs: tvm.tirx.call_packed("testing.echo", ins[0], ins[1],
outs[0]),
+ name="C",
+ )
+ D = te.compute(C.shape, lambda i, j: C[i, j] + 1.0, name="D")
+ return [A, B, D]
+
+
[email protected]_func(s_tir=True)
+def tir_extern_epilogue(var_A: T.handle, var_B: T.handle, D: T.Buffer((4, 2),
"float32")):
+ T.func_attr({"global_symbol": "main", "tirx.noalias": True})
+ A = T.match_buffer(var_A, (4, 3), offset_factor=1)
+ B = T.match_buffer(var_B, (3, 2), offset_factor=1)
+ C = T.sblock_alloc_buffer((4, 2), elem_offset=0, offset_factor=1)
+ with T.sblock("C"):
+ T.reads()
+ T.writes()
+ T.call_packed("testing.echo", A, B, C)
+ for i, j in T.grid(4, 2):
+ with T.sblock("D"):
+ vi, vj = T.axis.remap("SS", [i, j])
+ T.reads(C[vi, vj])
+ T.writes(D[vi, vj])
+ D[vi, vj] = C[vi, vj] + T.float32(1)
+
+
+def test_extern_epilogue():
+ _check_workload(te_extern_epilogue, tir_extern_epilogue)
+ func =
te.create_prim_func(te_extern_epilogue()).with_attr("global_symbol",
"extern_epilogue")
+ tvm.compile(func, target="llvm")
+
+
def te_reordered_matmul():
k = te.reduce_axis((0, 128), "k")
A = te.placeholder((128, 128), name="A")
diff --git a/tests/python/testing/test_testing.py
b/tests/python/testing/test_testing.py
deleted file mode 100644
index 373e78845b..0000000000
--- a/tests/python/testing/test_testing.py
+++ /dev/null
@@ -1,116 +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.
-# ruff: noqa: E731, F821, F841
-import numpy as np
-
-import tvm
-import tvm.testing
-
-
-def test_check_numerical_grads():
- # Functions and their derivatives
- functions = [
- lambda x: (x * x * x, 3 * x * x),
- lambda x: (x * x, 2 * x),
- lambda x: (np.abs(x), np.sign(x)),
- lambda x: (np.log(np.abs(x)), 1 / x),
- lambda x: (np.sqrt(np.abs(x)), np.sign(x) / (2 * np.sqrt(np.abs(x)))),
- lambda x: (1 / x, -1 / (x * x)),
- lambda x: (np.sign(np.sin(1 / x)), np.zeros_like(x)),
- lambda x: (x * np.sin(1 / x), np.sin(1 / x) - np.cos(1 / x) / x),
- lambda x: (np.sin(1 / x), -np.cos(1 / x) / (x * x)),
- lambda x: (np.tan(x), 1.0 / (np.cos(x) * np.cos(x))),
- ]
-
- np.random.seed(0)
-
- # Avoid values too close to 0 since singularities of our functions are
there
- min_x = 0.5
-
- for func in functions:
- x_input = np.random.uniform(min_x, 10, size=(3, 4))
-
- # We need a function returning a scalar, so sum the results
- func_forw = lambda x: np.sum(func(x)[0])
- grads = [func(x_input)[1]]
-
- tvm.testing.check_numerical_grads(func_forw, [x_input], grads)
-
- # Check functions with multiple arguments
- for f1 in functions:
- for f2 in functions:
- x_input = np.random.uniform(min_x, 10, size=(3, 4))
- y_input = np.random.uniform(min_x, 10, size=(3, 4))
-
- func_forw = lambda x, y: np.sum(f1(x)[0] + f2(y)[0])
- grads = [f1(x_input)[1], f2(y_input)[1]]
-
- tvm.testing.check_numerical_grads(func_forw, [x_input, y_input],
grads)
-
- # Same thing but with keyword arguments
- func_forw = lambda x, y: np.sum(f1(x)[0] + f2(y)[0])
- grads = {"x": f1(x_input)[1], "y": f2(y_input)[1]}
-
- tvm.testing.check_numerical_grads(func_forw, {"x": x_input, "y":
y_input}, grads)
-
- def _noise1(x, atol=1e-2, rtol=0.1):
- # We go in random direction using twice the original tolerance to be
sure this
- # results in an error
- sqrt_n = np.sqrt(float(np.prod(x.shape)))
- tol = 2 * (np.linalg.norm(x) * rtol + atol * sqrt_n)
- noise = np.random.normal(size=x.shape)
- noise = tol * noise / np.linalg.norm(noise)
- return x + noise
-
- def _noise2(x, atol=1e-2, rtol=0.1):
- # This noise affects just a single component
- sqrt_n = np.sqrt(float(np.prod(x.shape)))
- tol = 2 * (np.linalg.norm(x) * rtol + atol * sqrt_n)
- n = np.random.randint(np.prod(x.shape))
- noise = np.zeros_like(x)
- noise.reshape(-1)[n] = tol
- return x + noise
-
- # Add noise to gradients and check that the function throws
- for f1 in functions:
- for f2 in functions:
- x_input = np.random.uniform(min_x, 10, size=(3, 4))
- y_input = np.random.uniform(min_x, 10, size=(3, 4))
-
- func_forw = lambda x, y: np.sum(f1(x)[0] + f2(y)[0])
- grads = [_noise1(f1(x_input)[1]), _noise1(f2(y_input)[1])]
-
- try:
- tvm.testing.check_numerical_grads(func_forw, [x_input,
y_input], grads)
- except AssertionError as e:
- pass
- else:
- raise AssertionError("tvm.testing.check_numerical_grads didn't
raise an exception")
-
- func_forw = lambda x, y: np.sum(f1(x)[0] + f2(y)[0])
- grads = {"x": _noise2(f1(x_input)[1]), "y":
_noise2(f2(y_input)[1])}
-
- try:
- tvm.testing.check_numerical_grads(func_forw, {"x": x_input,
"y": y_input}, grads)
- except AssertionError as e:
- pass
- else:
- raise AssertionError("tvm.testing.check_numerical_grads didn't
raise an exception")
-
-
-if __name__ == "__main__":
- test_tvm.testing.check_numerical_grads()
diff --git a/tests/python/testing/test_tvm_testing_before_after.py
b/tests/python/testing/test_tvm_testing_before_after.py
deleted file mode 100644
index 7fb7cbbff0..0000000000
--- a/tests/python/testing/test_tvm_testing_before_after.py
+++ /dev/null
@@ -1,147 +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 tvm
-import tvm.testing
-from tvm.script import ir_module
-from tvm.script import tirx as T
-
-
-def test_before_after_prim_func():
- @T.prim_func(private=True, s_tir=True)
- def before():
- T.evaluate(0)
-
- expected = before
-
- mod = tvm.IRModule.from_expr(before)
- # Identity transform (no-op)
- mod = (lambda x: x)(mod)
- tvm.ir.assert_structural_equal(mod["main"], expected)
-
-
-def test_before_after_method():
- @T.prim_func(private=True, s_tir=True)
- def before():
- T.evaluate(0)
-
- expected = before
-
- mod = tvm.IRModule.from_expr(before)
- # Identity transform (no-op)
- mod = (lambda x: x)(mod)
- tvm.ir.assert_structural_equal(mod["main"], expected)
-
-
-def test_before_after_fixture():
- @T.prim_func(private=True, s_tir=True)
- def before():
- T.evaluate(0)
-
- expected = before
-
- mod = tvm.IRModule.from_expr(before)
- # Identity transform (no-op)
- mod = (lambda x: x)(mod)
- tvm.ir.assert_structural_equal(mod["main"], expected)
-
-
-def test_before_after_delayed_prim_func():
- @T.prim_func(private=True, s_tir=True)
- def before():
- T.evaluate(0)
-
- expected = before
-
- mod = tvm.IRModule.from_expr(before)
- # Identity transform (no-op)
- mod = (lambda x: x)(mod)
- tvm.ir.assert_structural_equal(mod["main"], expected)
-
-
-def test_before_after_parametrized_fixture():
- """Test with different buffer sizes"""
- for n in [1, 8, 16]:
-
- @T.prim_func(private=True, s_tir=True)
- def before(A: T.Buffer(n, "float32")):
- for i in T.serial(n):
- A[i] = 0.0
-
- expected = before
-
- mod = tvm.IRModule.from_expr(before)
- # Identity transform (no-op)
- mod = (lambda x: x)(mod)
- tvm.ir.assert_structural_equal(mod["main"], expected)
-
-
-def test_before_after_ir_module():
- """The preferred form for writing TIR unit tests
-
- All evaluation is done at test-time, with the minimal amount of
- additional lines.
- """
-
- @ir_module
- class before:
- @T.prim_func(private=True, s_tir=True)
- def func_A(A: T.Buffer(16, "float32")):
- for i in T.serial(16):
- A[i] = 0.0
-
- @T.prim_func(private=True, s_tir=True)
- def func_B(A: T.Buffer(16, "int32")):
- for i in T.serial(16):
- A[i] = 42
-
- expected = before
-
- # Identity transform (no-op)
- mod = (lambda x: x)(before)
- tvm.ir.assert_structural_equal(mod, expected)
-
-
-def test_before_after_ir_module_explicit_fixture():
- """Like test_before_after_ir_module, but with an explicit fixture
-
- If the IRModule depends on additional fixtures, this form can be
- used.
- """
-
- @ir_module
- class before:
- @T.prim_func(private=True, s_tir=True)
- def func_A(A: T.Buffer(16, "float32")):
- for i in T.serial(16):
- A[i] = 0.0
-
- @T.prim_func(private=True, s_tir=True)
- def func_B(A: T.Buffer(16, "int32")):
- for i in T.serial(16):
- A[i] = 42
-
- expected = before
-
- # Identity transform (no-op)
- mod = (lambda x: x)(before)
- tvm.ir.assert_structural_equal(mod, expected)
-
-
-if __name__ == "__main__":
- tvm.testing.main()
diff --git a/tests/python/testing/test_tvm_testing_features.py
b/tests/python/testing/test_tvm_testing_features.py
deleted file mode 100644
index 2f4c798947..0000000000
--- a/tests/python/testing/test_tvm_testing_features.py
+++ /dev/null
@@ -1,192 +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.
-# ruff: noqa: RUF012
-
-import os
-
-import pytest
-
-import tvm.testing
-
-pytestmark = pytest.mark.xdist_group(name="tvm-testing-features")
-
-# This file tests features in tvm.testing, such as verifying that
-# cached fixtures are run an appropriate number of times. As a
-# result, the order of the tests is important. Use of --last-failed
-# or --failed-first while debugging this file is not advised. If
-# these tests are distributed/parallelized using pytest-xdist or
-# similar, all tests in this file should run sequentially on the same
-# node. (See https://stackoverflow.com/a/59504228)
-
-
-class TestParameter:
- param1_vals = [1, 2, 3]
- param2_vals = ["a", "b", "c"]
-
- independent_usages = 0
- param1 = tvm.testing.parameter(*param1_vals)
- param2 = tvm.testing.parameter(*param2_vals)
-
- def test_using_independent(self, param1, param2):
- type(self).independent_usages += 1
-
- def test_independent(self):
- assert self.independent_usages == len(self.param1_vals) *
len(self.param2_vals)
-
-
-class TestFixtureCaching:
- param1_vals = [1, 2, 3]
- param2_vals = ["a", "b", "c"]
-
- param1 = tvm.testing.parameter(*param1_vals)
- param2 = tvm.testing.parameter(*param2_vals)
-
- uncached_calls = 0
- cached_calls = 0
-
- @tvm.testing.fixture
- def uncached_fixture(self, param1):
- type(self).uncached_calls += 1
- return 2 * param1
-
- def test_use_uncached(self, param1, param2, uncached_fixture):
- assert 2 * param1 == uncached_fixture
-
- def test_uncached_count(self):
- assert self.uncached_calls == len(self.param1_vals) *
len(self.param2_vals)
-
- @tvm.testing.fixture(cache_return_value=True)
- def cached_fixture(self, param1):
- type(self).cached_calls += 1
- return 3 * param1
-
- def test_use_cached(self, param1, param2, cached_fixture):
- assert 3 * param1 == cached_fixture
-
- def test_cached_count(self):
- cache_disabled = bool(int(os.environ.get("TVM_TEST_DISABLE_CACHE",
"0")))
- if cache_disabled:
- assert self.cached_calls == len(self.param1_vals) *
len(self.param2_vals)
- else:
- assert self.cached_calls == len(self.param1_vals)
-
-
-class TestCachedFixtureIsCopy:
- param = tvm.testing.parameter(1, 2, 3, 4)
-
- @tvm.testing.fixture(cache_return_value=True)
- def cached_mutable_fixture(self):
- return {"val": 0}
-
- def test_modifies_fixture(self, param, cached_mutable_fixture):
- assert cached_mutable_fixture["val"] == 0
-
- # The tests should receive a copy of the fixture value. If
- # the test receives the original and not a copy, then this
- # will cause the next parametrization to fail.
- cached_mutable_fixture["val"] = param
-
-
-class TestBrokenFixture:
- # Tests that use a fixture that throws an exception fail, and are
- # marked as setup failures. The tests themselves are never run.
- # This behavior should be the same whether or not the fixture
- # results are cached.
-
- num_uses_broken_uncached_fixture = 0
- num_uses_broken_cached_fixture = 0
-
- @tvm.testing.fixture
- def broken_uncached_fixture(self):
- raise RuntimeError("Intentionally broken fixture")
-
- @pytest.mark.xfail(True, reason="Broken fixtures should result in a
failing setup", strict=True)
- def test_uses_broken_uncached_fixture(self, broken_uncached_fixture):
- type(self).num_uses_broken_fixture += 1
-
- def test_num_uses_uncached(self):
- assert self.num_uses_broken_uncached_fixture == 0
-
- @tvm.testing.fixture(cache_return_value=True)
- def broken_cached_fixture(self):
- raise RuntimeError("Intentionally broken fixture")
-
- @pytest.mark.xfail(True, reason="Broken fixtures should result in a
failing setup", strict=True)
- def test_uses_broken_cached_fixture(self, broken_cached_fixture):
- type(self).num_uses_broken_cached_fixture += 1
-
- def test_num_uses_cached(self):
- assert self.num_uses_broken_cached_fixture == 0
-
-
[email protected](
- bool(int(os.environ.get("TVM_TEST_DISABLE_CACHE", "0"))),
- reason="Cannot test cache behavior while caching is disabled",
-)
-class TestCacheableTypes:
- class EmptyClass:
- pass
-
- @tvm.testing.fixture(cache_return_value=True)
- def uncacheable_fixture(self):
- return self.EmptyClass()
-
- def test_uses_uncacheable(self, request):
- # Normally the num_tests_use_this_fixture would be set before
- # anything runs. For this test case only, because we are
- # delaying the use of the fixture, we need to manually
- # increment it.
- self.uncacheable_fixture.num_tests_use_this_fixture[0] += 1
- with pytest.raises(TypeError):
- request.getfixturevalue("uncacheable_fixture")
-
- class ImplementsReduce:
- def __reduce__(self):
- return super().__reduce__()
-
- @tvm.testing.fixture(cache_return_value=True)
- def fixture_with_reduce(self):
- return self.ImplementsReduce()
-
- def test_uses_reduce(self, fixture_with_reduce):
- pass
-
- class ImplementsDeepcopy:
- def __deepcopy__(self, memo):
- return type(self)()
-
- @tvm.testing.fixture(cache_return_value=True)
- def fixture_with_deepcopy(self):
- return self.ImplementsDeepcopy()
-
- def test_uses_deepcopy(self, fixture_with_deepcopy):
- pass
-
-
-class TestPytestCache:
- param = tvm.testing.parameter(1, 2, 3)
-
- @pytest.fixture(scope="class")
- def cached_fixture(self, param):
- return param * param
-
- def test_uses_cached_fixture(self, param, cached_fixture):
- assert cached_fixture == param * param
-
-
-if __name__ == "__main__":
- tvm.testing.main()
diff --git a/tests/python/testing/test_type_annotation_checker.py
b/tests/python/testing/test_type_annotation_checker.py
deleted file mode 100644
index b5d2afcb92..0000000000
--- a/tests/python/testing/test_type_annotation_checker.py
+++ /dev/null
@@ -1,227 +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.
-# ruff: noqa: F401
-"""Test type checker based on python's type annotations"""
-
-import sys
-from collections.abc import Callable
-from typing import Union
-
-import _pytest
-import pytest
-
-import tvm
-from tvm.s_tir.schedule._type_checker import type_checked
-
-
-def int_func(x: int) -> int:
- return 2 * x
-
-
-def str_func(x: str) -> str:
- return 2 * x
-
-
-test_cases = [
- {
- "type_annotation": int,
- "positive_cases": [5],
- "negative_cases": ["5"],
- },
- {
- "type_annotation": list[int],
- "positive_cases": [
- [5],
- [],
- # Tuples are allowed to be used as lists, because both are
- # represented in FFI as tvm::Array.
- (1, 2, 3),
- ],
- "negative_cases": [
- None,
- 5,
- ["5"],
- ],
- },
- {
- "type_annotation": dict[str, int],
- "positive_cases": [
- {"key1": 0, "key2": 1, "key3": -1},
- ],
- "negative_cases": [None, [1], {1: "1"}],
- },
- {
- "type_annotation": tuple[int],
- "positive_cases": [
- (5,),
- ],
- "negative_cases": [
- None,
- (1, 2, 3),
- [1],
- 5,
- ["5"],
- ],
- },
- {
- "type_annotation": tuple[str, int],
- "positive_cases": [
- ("x", 5),
- ],
- "negative_cases": [
- 42,
- ("x", 5, 6),
- ("x", 5, "y"),
- ("x", 5.0),
- (None, 5),
- ],
- },
- {
- "type_annotation": str | int,
- "positive_cases": [
- "x",
- 5,
- ],
- "negative_cases": [
- 5.0,
- ("x", 5, 6),
- None,
- ],
- },
- {
- "type_annotation": Callable,
- "positive_cases": [str_func, int_func],
- "negative_cases": [
- None,
- "x",
- 42,
- ],
- },
- {
- "type_annotation": Callable[[int], int],
- "positive_cases": [int_func],
- "negative_cases": [
- None,
- "x",
- 42,
- pytest.param(
- str_func,
- marks=pytest.mark.xfail(
- reason="Signature of Callable arguments not currently
checked"
- ),
- ),
- ],
- },
-]
-
-
-def make_parametrization(type_annotation, case):
- if isinstance(case, _pytest.mark.structures.ParameterSet):
- marks = case.marks
- (case,) = case.values
- else:
- marks = []
-
- try:
- annotation_name = type_annotation.__name__
- except AttributeError:
- annotation_name = str(type_annotation).replace("typing.", "")
-
- if hasattr(case, "__name__"):
- case_name = case.__name__
- else:
- case_name = str(case)
-
- name = f"{annotation_name}, {case_name}"
-
- return pytest.param(type_annotation, case, marks=marks, id=name)
-
-
-positive_cases = [
- make_parametrization(config["type_annotation"], case)
- for config in test_cases
- for case in config["positive_cases"]
-]
-
-negative_cases = [
- make_parametrization(config["type_annotation"], case)
- for config in test_cases
- for case in config["negative_cases"]
-]
-
-
[email protected](
- ["type_annotation", "case"],
- positive_cases,
-)
-def test_matches_type(type_annotation, case):
- @type_checked
- def func(_: type_annotation):
- pass
-
- func(case)
-
-
[email protected](
- ["type_annotation", "case"],
- negative_cases,
-)
-def test_not_matches(type_annotation, case):
- @type_checked
- def func(_: type_annotation):
- pass
-
- with pytest.raises(TypeError):
- func(case)
-
-
[email protected](
- ["type_annotation", "expected_key", "expected_subtypes"],
- [
- pytest.param(str | int, "union", [str, int], id="str | int"),
- pytest.param(list[str], "list", [str], id="List[str]"),
- pytest.param(dict[str, int], "dict", [str, int], id="Dict[str, int]"),
- pytest.param(tuple[str, int], "tuple", (str, int), id="Tuple[str,
int]"),
- pytest.param(
- list[str] | dict[str, int],
- "union",
- [list[str], dict[str, int]],
- id="Union[List[str], Dict[str, int]]",
- ),
- ],
-)
-def test_subscripted_generics(type_annotation, expected_key,
expected_subtypes):
- """Test that _dispatcher correctly handles subscripted generics in Python
3.14+.
-
- In Python 3.14, Union and other generic types have a different internal
representation.
- This test ensures that the dispatcher correctly identifies these types.
- """
- from tvm.s_tir.schedule._type_checker import _dispatcher
-
- key, subtypes = _dispatcher(type_annotation)
- assert key == expected_key, f"Expected '{expected_key}' but got '{key}'"
-
- if isinstance(expected_subtypes, tuple):
- assert tuple(subtypes) == expected_subtypes, (
- f"Expected {expected_subtypes} but got {subtypes}"
- )
- else:
- assert subtypes == expected_subtypes, f"Expected {expected_subtypes}
but got {subtypes}"
-
-
-if __name__ == "__main__":
- tvm.testing.main()
diff --git
a/tests/python/tirx-transform/test_tir_transform_lower_tvm_builtin.py
b/tests/python/tirx-transform/test_tir_transform_lower_tvm_builtin.py
index df3e8d529c..d58d993d2a 100644
--- a/tests/python/tirx-transform/test_tir_transform_lower_tvm_builtin.py
+++ b/tests/python/tirx-transform/test_tir_transform_lower_tvm_builtin.py
@@ -113,6 +113,36 @@ def test_lower_call_packed():
tvm.ir.assert_structural_equal(After, Expected)
[email protected](not env.has_llvm(), reason="need llvm")
+def test_lower_call_packed_raw_string():
+ @I.ir_module
+ class Before:
+ @T.prim_func(s_tir=True)
+ def main():
+ T.func_attr({"target": tvm.target.Target("llvm")})
+ T.call_packed("testing.echo", "payload")
+
+ @I.ir_module
+ class Expected:
+ @T.prim_func(s_tir=True)
+ def main():
+ T.func_attr({"target": tvm.target.Target("llvm")})
+ stack_ffi_any: T.let[T.handle] = T.tvm_stack_alloca("tvm_ffi_any",
2)
+ T.tvm_struct_set(stack_ffi_any, 0, 13, 8)
+ T.tvm_struct_set(stack_ffi_any, 0, 14, 0)
+ T.tvm_struct_set(stack_ffi_any, 0, 15,
T.reinterpret(T.handle().ty, "payload"))
+ T.tvm_struct_set(stack_ffi_any, 1, 13, 0)
+ T.tvm_struct_set(stack_ffi_any, 1, 14, 0)
+ T.tvm_struct_set(stack_ffi_any, 1, 15, T.int64(0))
+ T.call_packed_lowered("testing.echo", stack_ffi_any, 0, 1)
+
+ After = tvm.tirx.transform.LowerTVMBuiltin()(Before)
+ tvm.ir.assert_structural_equal(After, Expected)
+
+ # The typed pointer is required by the LLVM TVMFFIAny lowering.
+ tvm.compile(Before, target="llvm")
+
+
@pytest.mark.skipif(not env.has_llvm(), reason="need llvm")
def test_call_packed_return_non_i32():
# This call packed that return non i32 types
diff --git a/tests/python/tirx/conftest.py b/tests/python/tirx/conftest.py
index fb8ba62f4f..2653a1f05a 100644
--- a/tests/python/tirx/conftest.py
+++ b/tests/python/tirx/conftest.py
@@ -25,6 +25,8 @@ real sm_100a device so it skips cleanly where the hardware is
absent and runs
in full where it is present.
"""
+from pathlib import Path
+
import pytest
from tvm.testing import env
@@ -33,8 +35,17 @@ from tvm.testing import env
def pytest_collection_modifyitems(config, items):
if env.has_cuda_compute(10):
return
+ suite_root = Path(__file__).resolve().parent
skip = pytest.mark.skip(
reason="tirx suite requires a CUDA compute capability 10.0 (sm_100a)
device"
)
for item in items:
- item.add_marker(skip)
+ path = getattr(item, "path", None)
+ if path is None:
+ continue
+ try:
+ path = Path(path).resolve()
+ except TypeError:
+ continue
+ if path.is_relative_to(suite_root):
+ item.add_marker(skip)
diff --git a/tests/python/tvmscript/test_tvmscript_roundtrip.py
b/tests/python/tvmscript/test_tvmscript_roundtrip.py
index 10828ddf25..08bf42decf 100644
--- a/tests/python/tvmscript/test_tvmscript_roundtrip.py
+++ b/tests/python/tvmscript/test_tvmscript_roundtrip.py
@@ -3292,15 +3292,11 @@ def relax_symbolic_var():
def relax_float_symbolic_var():
- """Relax symbolic variables may hold any dtype"""
+ """Relax scalar variables may use any dtype."""
@R.function
- def func(A: R.Tensor(["N"], "float16"), _: R.Prim(value="threshold")):
- N = T.int64()
- threshold = T.float16()
-
- B = A >= R.prim_value(threshold / T.cast(N, "float16"))
- return B
+ def func(value: R.Prim("float16")):
+ return value
return func
diff --git a/tests/scripts/ci.py b/tests/scripts/ci.py
index bea70e0b52..fa027c150a 100755
--- a/tests/scripts/ci.py
+++ b/tests/scripts/ci.py
@@ -591,7 +591,6 @@ generated = [
[
"./tests/scripts/task_java_unittest.sh",
"./tests/scripts/task_python_unittest_gpuonly.sh",
- "./tests/scripts/task_python_integration_gpuonly.sh",
],
),
},
@@ -601,10 +600,6 @@ generated = [
help="Run CPU build and test(s)",
options={
"cpp": CPP_UNITTEST,
- "integration": (
- "run integration tests",
- ["./tests/scripts/task_python_integration.sh"],
- ),
"unittest": (
"run unit tests",
[
diff --git a/tests/scripts/task_clear_pytest.sh
b/tests/scripts/task_clear_pytest.sh
deleted file mode 100755
index 430425791c..0000000000
--- a/tests/scripts/task_clear_pytest.sh
+++ /dev/null
@@ -1,22 +0,0 @@
-#!/usr/bin/env bash
-# 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.
-
-set -euo pipefail
-
-# PR jobs evaluate their Jenkinsfile from the trusted base branch. Keep this
-# inert entry point until the Jenkinsfile update has landed on the base branch.
diff --git a/tests/scripts/task_python_integration.sh
b/tests/scripts/task_python_integration.sh
deleted file mode 100755
index 7dd3c9b578..0000000000
--- a/tests/scripts/task_python_integration.sh
+++ /dev/null
@@ -1,32 +0,0 @@
-#!/usr/bin/env bash
-# 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.
-
-set -euxo pipefail
-
-export PYTHONPATH="$(pwd)/python"
-export LD_LIBRARY_PATH="build:${LD_LIBRARY_PATH:-}"
-
-# to avoid CI CPU thread throttling.
-export TVM_BIND_THREADS=0
-export TVM_NUM_THREADS=2
-
-# cleanup pycache
-find . -type f -path "*.pyc" | xargs rm -f
-
-# setup tvm-ffi into python folder
-uv pip install -v --target=python ./3rdparty/tvm-ffi/
diff --git a/tests/scripts/task_python_integration_gpuonly.sh
b/tests/scripts/task_python_integration_gpuonly.sh
deleted file mode 100755
index 9b825b9182..0000000000
--- a/tests/scripts/task_python_integration_gpuonly.sh
+++ /dev/null
@@ -1,27 +0,0 @@
-#!/usr/bin/env bash
-# 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.
-
-set -exo pipefail
-
-export
TVM_TEST_TARGETS='cuda;opencl;metal;rocm;nvptx;{"kind":"opencl","device":"mali,adreno"}'
-# Every GPU test carries the `gpu` marker; the specific backend is gated by
skipif.
-export PYTEST_ADDOPTS="-m gpu $PYTEST_ADDOPTS"
-export TVM_RELAY_TEST_TARGETS="cuda"
-export TVM_INTEGRATION_GPU_ONLY=1
-
-./tests/scripts/task_python_integration.sh
diff --git a/tests/scripts/task_python_unittest.sh
b/tests/scripts/task_python_unittest.sh
index 0728743fb5..524a9f2a2e 100755
--- a/tests/scripts/task_python_unittest.sh
+++ b/tests/scripts/task_python_unittest.sh
@@ -19,55 +19,9 @@
set -euxo pipefail
export PYTHONPATH="$(pwd)/python"
-export PYTEST_ADDOPTS="-s -vv ${CI_PYTEST_ADD_OPTIONS:-} ${PYTEST_ADDOPTS:-}"
-
-# cleanup pycache
-find . -type f -path "*.pyc" | xargs rm -f
+export PYTEST_ADDOPTS="${CI_PYTEST_ADD_OPTIONS:-} ${PYTEST_ADDOPTS:-}"
# setup tvm-ffi into python folder
uv pip install -v --target=python ./3rdparty/tvm-ffi/
-# First run the minimal platform test. A GPU-only run can select no tests
here.
-if [ ! -d tests/python/all-platform-minimal-test ]; then
- echo "Missing pytest target: tests/python/all-platform-minimal-test" >&2
- exit 1
-fi
-python3 -m pytest -n auto tests/python/all-platform-minimal-test || [ "$?" -eq
5 ]
-
-# Then run all unit tests.
-TEST_FILES=(
- "arith"
- "ci"
- "codegen"
- "driver"
- "ir"
- "runtime"
- "target"
- "te"
- "testing"
- "s_tir/base"
- "s_tir/schedule"
- "s_tir/dlight"
- "s_tir/analysis"
- "s_tir/meta_schedule"
- "s_tir/transform"
- "tirx-analysis"
- "tirx-base"
- "tirx-transform"
- "tirx"
- "tvmscript"
- "relax"
-)
-
-PYTEST_TARGETS=()
-for TEST_FILE in "${TEST_FILES[@]}"; do
- TEST_PATH="tests/python/${TEST_FILE}"
- if [ ! -e "${TEST_PATH}" ]; then
- echo "Missing pytest target: ${TEST_PATH}" >&2
- exit 1
- fi
- PYTEST_TARGETS+=("${TEST_PATH}")
-done
-
-# Do not mask pytest's exit 5: an unexpectedly empty broad suite must fail CI.
-python3 -m pytest -n auto --dist=loadgroup "${PYTEST_TARGETS[@]}"
+python3 -m pytest -vvs -n auto -m "${TVM_TEST_MARKER:-not gpu}" tests/python
diff --git a/tests/scripts/task_python_unittest_gpuonly.sh
b/tests/scripts/task_python_unittest_gpuonly.sh
index 8194b23ac4..ca1ae86f73 100755
--- a/tests/scripts/task_python_unittest_gpuonly.sh
+++ b/tests/scripts/task_python_unittest_gpuonly.sh
@@ -19,7 +19,7 @@
set -euxo pipefail
# Every GPU test carries the `gpu` marker; the specific backend is gated by
skipif.
-export PYTEST_ADDOPTS="-m gpu ${PYTEST_ADDOPTS:-}"
+export TVM_TEST_MARKER="gpu"
# Test most of the enabled runtimes here.
# TODO: disabled opencl tests due to segmentation fault.
@@ -32,10 +32,7 @@ export TVM_TEST_TARGETS='cuda;metal;rocm;nvptx'
export TVM_TEST_TARGETS='{"kind":"vulkan","from_device":0}'
export PYTHONPATH="$(pwd)/python"
-export PYTEST_ADDOPTS="-s -vv ${CI_PYTEST_ADD_OPTIONS:-} ${PYTEST_ADDOPTS:-}"
+export PYTEST_ADDOPTS="${CI_PYTEST_ADD_OPTIONS:-} ${PYTEST_ADDOPTS:-}"
-if [ ! -f tests/python/codegen/test_target_codegen_vulkan.py ]; then
- echo "Missing pytest target:
tests/python/codegen/test_target_codegen_vulkan.py" >&2
- exit 1
-fi
-python3 -m pytest -n auto tests/python/codegen/test_target_codegen_vulkan.py
+python3 -m pytest -vvs -n auto -m "${TVM_TEST_MARKER}" \
+ tests/python/codegen/test_target_codegen_vulkan.py
