Add tvm.testing.run_with_gpu_lock backed by the existing
tvm_ffi.utils.FileLock. Migrate live local GPU tests to acquire the
machine-local lock around device execution, synchronization, host
transfer, and checks while leaving target construction and compilation
outside the critical section.
Replace the custom xdist scheduler with standard xdist_group placement
for the order-dependent test family. RPC tests retain dynamic port
allocation and per-test process isolation rather than gaining a broad
category lock.
Drop accumulated dead code in the test-support package: helpers with
zero call sites, unused capability probes, dead FFI re-exports, and
orphaned pytest plumbing. Verified by repo-wide grep that nothing
references any of these.
This pr modernizes test gating. It replaces the heavy
`tvm.testing.Feature` machinery with a thin `tvm.testing.env` module of
`has_*()` capability probes, used via standard pytest.mark + skipif. And
markers move to `pyproject.toml`
[REFACTOR] Phase out getattr based attribute handling
This PR phases out getattar based attribute handling as they are slower
and introduces extra code path.
This does mean that if an Object is not explicitly registered
in python side, we will no longer be able to access the field by name.
Likely this is also desirable as we would like to enable faster use that
updates the python end and do not rely on these behavior.
This PR modernizes the FFI foundation of the project and introduce
a new minimal and lightweight module [tvm ffi](https://github.com/apache/tvm/tree/refactor-s3/ffi)
based on our lessons in the past few years. It implements a modern
version of the [Unified Packed and Object RFC](https://github.com/apache/tvm-rfcs/blob/main/rfcs/0097-unify-packed-and-object.md)
that unifies the packed function call and object systems.
Summary of the change:
- A dedicated clean Any/AnyView that can store strong and weak
references of items
- Function(previously PackedFunc) system built on top of the Any/AnyView
- A minimal C API that backs the overall calls. We are stabilizing the
API with a goal to bring clean, stable FFI conventions for both compiled
and registered code
- A rewrite of core python binding and generated code based on the module
- Update existing code and test cases to the new module
- Latest dlpack support
The new module brings many benefits thanks to the cleaner design,
to name a few:
- Any can support both POD types(int) and object types.
- Containers (e.g. Array) can now also contain Any value, e.g. now
`Array<int>` is supported, no need for boxed types
- Error handling now upgrades to object-based, allowing cleaner
traceback across languages
- Map now preserves insertion orders
- Path toward isolated stabilize minimum core ABI/API foundation module
- Type traits based design that cleanly defines how values interact
with Any system
- Automatic conversion of different types based on traits if needed
Because FFI upgrade is at heart of the project, the change touches every
component of the system. Importantly, this is an upgrade of the ABI so the
change is not backward compatible. The code compiled under the old
FFI won't work under the new one. We did provide example ABI translation
(e.g. LegacyTVMArgValueToFFIAny) functions for compatibility.
The PR tries to leave files in their old places while creating redirections.
The goal is to have the first milestone landed and infrastructure in place,
so we can do further refactors to complete features and cleanup legacy code
as trackable PRs. As of now, python binding and compiled code are under the
new convention while RPC and some other bindings still relies on legacy ABI
translation. We will work on upgrades in the coming PRs, including areas such
as reflection, phasing out legacy redirections etc.
This PR starts the step 0 to phase out relay from the current
development main branch. This PR focuses on the python
components of relay, autotvm, auto_scheduler. To make the change
manageable, we will also do followup steps on te.Schedule and
c++ components in followup PRs.
To continue support community members who depends on
legacy flows, the [v0.19.0](https://github.com/apache/tvm/tree/v0.19.0)
branch will continue contain these components.
As noted in [discussion on phasing out legacy components](https://discuss.tvm.apache.org/t/phasing-out-legacy-components/17703/30),
this would help us to do two purposes:
- By removing outdated or redundant elements, we can significantly
reduce complexity and improve maintainability.
- Unify our focus: Concentrating our efforts on the new unity flow
will allow for more efficient development and innovation.
It is also a good opportunity for us to revisit and reduce CI time.
The past relay legacy flow contains a lot of end to end tests that
requires hardware resources to run and causing long CI time.
Moving onwards, we can focus more on unit-tests that focuses
on structural equality and runs within seconds, while be mindful
about tests that requires hardware resources (by restricting them
to specific folders and CI nightly in some cases).
---
Co-authored-by: Siyuan Feng <hzfengsy@sjtu.edu.cn>
Disco worker originally automatically import `tvm.testing.disco` for
convenient unittesting. However, `tvm.testing` is a special subpackage
that introduces many unnecessary dependencies, for example, pytest. This
PR removes such dependencies by directly moving the testing function
registration logic to the entry file.
This PR introduces `ProcessSession`, a new session implementation based
on multi-processing.
`ProcessSession` shares exactly the same communication protocol with
`ThreadedSession`, but all workers except for worker 0 are launched in a
separate process than thread. Workers communicate with the controller
via pipe provided by the OS, rather than SPSC message queue between
threads.
In our implementation, Python's `subproces.popen` is used to create
subprocesses, and the Python executable, or more specifically,
`sys.executable` calls into `tvm.exec.disco_worker` as the entrypoint.
Besides the launching logic that is only executed once in the very
beginning, the rest of the implementation resides in a C++-only
environment, including reads/writes to pipe file descriptors,
serialization and deserialization of messages, worker interpretation of
each message, etc.
Detailed engineering elements included in this PR:
- Refactors the MinRPC-based communication protocol out to be shared by
`ProcessSession` and `ThreadedSession` as `protocol.h`;
- Refactors a controller-side worker thread into `DiscoWorkerThread`,
which is shared by both session implementation to launch worker-0;
- Added two instructions `kDebugGetFromRemote` and `kDebugSetRegister`,
which are used to communicate with workers other than worker-0 in
debug mode;
- Introduces multi-processing infra including: `tvm.exec.disco_worker`
serving as the entrypoint that launches workers, and
`tvm/runtime/disco/process_pool.py` that exposes APIs to launch worker
processes. `tvm.exec.disco_worker` calls into a global function
`runtime.disco.WorkerProcess` that executes the worker main loop in
pure C++;
- Introduces `src/support/process_id.h` that provides cross-platform pid
and tid printing utilities;
- Refactors Disco's NCCL integration that get rids of initialized-once
global NCCL context, and switches to broadcasting `ncclUniqueId` from
controller to all workers, and then create NCCL communicators in each
worker thread/process accordingly. This is a thread/process-agnostic
way of using NCCL.
This PR introduces `tvm.testing.rpc_run`, a utility method that allows a
`runtime.Module` to run on a remote device via TVM RPC.
Example:
```python
import numpy as np
import tvm
from tvm.script import tir as T
from tvm.testnig import rpc_run
@T.prim_func
def cuda_kernel(
A: T.Buffer((128,), "float32"),
B: T.Buffer((128,), "float32"),
):
for bx in T.thread_binding(4, thread="blockIdx.x"):
for tx in T.thread_binding(32, thread="threadIdx.x"):
x = bx * 32 + tx
B[x] = A[x] + 1.0
def main():
np_a = np.random.randn(128).astype("float32")
np_b = np_a + 1.0
rt_mod = tvm.build(cuda_kernel, target="nvidia/geforce-rtx-3090-ti")
tvm_a, tvm_b = rpc_run(
rt_mod,
"cuda",
[np_a, np_b],
)
assert np.allclose(tvm_b, np_b)
```
Result:
```
Execution time summary:
mean (ms) median (ms) max (ms) min (ms) std (ms)
0.0023 0.0023 0.0023 0.0023 0.0000
```
* [UnitTest] Added ids argument to tvm.testing.parameters
This matches the usage in `tvm.testing.parameter`, and allows for
parameter sets to be referred to by a single name.
* [Pytest] Fixed ordering issue of tvm.testing.parametrize_targets and known_failing_targets
If an explicit list of targets is given, then the
`known_failing_targets` decorator would fail to apply. This commit
resolves the issue, and cleans up all target-specific marks to apply
in `tvm.testing.plugin._add_target_specific_marks`.
* [UnitTest][Vulkan] Runnable relay unit tests on Vulkan
This commit allows the relay test suite to be run targeting Vulkan with
`TVM_TEST_TARGETS="vulkan -from_device=0" pytest tests/python/relay`. All
tests that require a specific environment are skipped if that environment
isn't present. All tests that are known to fail when running on Vulkan
are marked as expected failure, and will be tracked in
https://github.com/apache/tvm/issues/8903.
- Failures during code generation
- Type mismatches, boolean vs int8
- tests/python/relay/test_any.py::test_any_reduce
- tests/python/relay/test_op_level3.py::test_sparse_reshape
- tests/python/relay/test_op_level4.py::test_reduce_functions
- tests/python/relay/test_vm.py::test_cond
- tests/python/relay/test_vm.py::test_simple_if
- Incorrect strategy selection, picks NCHWc implemenation for NHWC layout
- tests/python/relay/test_op_level2.py::test_conv2d_run
- Unresolved CallNode operation
- tests/python/relay/test_op_level1.py::test_unary_op[erf/tan/atan]
- tests/python/relay/test_op_level3.py::test_scatter_add
- tests/python/relay/test_op_level3.py::test_segment_sum
- Generates 64-bit calls to GLSL that have only 16-/32-bit support
- tests/python/relay/test_op_grad_level1.py::test_log_softmax_grad
- tests/python/relay/test_op_grad_level1.py::test_softmax_grad
- tests/python/relay/test_op_grad_level1.py::test_unary_op
- tests/python/relay/test_op_grad_level10.py::test_cross_entropy_grad
- Codegen raises error for variable size
- tests/python/relay/test_any.py::test_any_batch_matmul
- tests/python/relay/test_any.py::test_any_conv2d_NCHWc
- tests/python/relay/test_any.py::test_any_dense
- Failures when running
- Numeric differences (observed on GTX 1650 with NVIDIA driver)
- tests/python/relay/test_op_level3.py::test_take
- tests/python/relay/test_op_level5.py::TestCropAndResize
- tests/python/relay/test_op_level5.py::TestResize1D
- tests/python/relay/test_op_level5.py::TestResize2D
- Adds the decorator `tvm.testing.requires_ethosn`
- Marks all tests in `tests/python/contrib/test_ethosn` as requiring
ethosn instead of directly checking `ethosn_available()`. This way,
they show up as skipped rather than passing.
- Marks test_compile_tflite_module_with_external_codegen as requiring
ethosn.
* [UnitTests][CoreML] Marked test_annotate as a known failure.
The unit tests in `test_coreml_codegen.py` haven't run in the CI
lately, so this test wasn't caught before. (See tracking issue
- Added `pytest.mark.xfail` mark to `test_annotate`.
- Added `tvm.testing.requires_package` decorator, which can mark tests
as requiring a specific python package to be available. Switched
from `pytest.importorskip('coremltools')` to
`requires_package('coremltools')` in `test_coreml_codegen.py` so
that all tests would explicitly show up as skipped in the report.
- Added `uses_gpu` tag to all tests in `test_coreml_codegen.py`, since
only ci_gpu has coremltools installed. In the future, if the ci_cpu
image has coremltools installed, this mark can be removed.
* [Pytest][TensorRT] Mark the TensorRT tests with tvm.testing.requires_cuda
Previously, the tests had an early bailout if tensorrt was disabled,
or if there was no cuda device present. However, the tests were not
marked with `pytest.mark.gpu` and so they didn't run during
`task_python_integration_gpuonly.sh`. This commit adds the
`requires_cuda` mark, and maintains the same behavior of testing the
tensorrt compilation steps if compilation is enabled, and running the
results if tensorrt is enabled.
In addition, some of the tests result in failures when run. These
have been marked with `pytest.mark.xfail`, and are being tracked in
issue #8901.
* [UnitTests] Expose TVM pytest helpers as plugin
Previously, pytest helper utilities such as automatic parametrization
of `target`/`dev`, or `tvm.testing.parameter` were only available for
tests within the `${TVM_HOME}/tests` directory. This PR extracts the
helper utilities into an importable plugin, which can be used in
external tests (e.g. one-off debugging).
* [UnitTests] Refactor the plugin-specific logic out into plugin.py.
* [UnitTests] Moved marker definition out to global variable.
Tianqi Chen