Repository navigation
[RELEASE] Backport main to prepare v0.25.0.rc1 - #19774
Conversation
|
NOTE ON PR MERGE: we need "Rebase and merge" instead of "Squash and merge" to keep the linear history. |
There was a problem hiding this comment.
Code Review
This pull request refactors TVM's error handling, FFI, and library loading mechanisms, notably replacing the legacy C++ diagnostic reporting system with a pass-time error enrichment mechanism that renders underlined TVMScript locations. It also cleans up library loading, transitions TVMError to standard RuntimeError, and updates Python dependencies to lazily import optional packages. The review feedback is highly constructive, pointing out critical issues such as a missing os import in emcc.py, a Python 3.9 compatibility issue with entry_points() in _autoload_backends.py, and potential AttributeErrors in the PyTorch FX translator when accessing struct_info.dtype directly. Additionally, it suggests minor optimizations for AST traversal in the GPU fallback schedule and more robust socket connections in the Hexagon pytest plugin.
Important
The consumer version of Gemini Code Assist on GitHub is being sunset. Starting June 18, 2026, new organization installations will be blocked, and all code review activity will officially cease on July 17, 2026.
For more details on the timeline and next steps, please review the Help Documentation.
1a4f891 to
6615ea2
Compare
|
@tvm-bot rerun |
1 similar comment
|
@tvm-bot rerun |
apache#19679) ### Motivation `torch.logical_and` accepts input tensors of any dtype (treating any nonzero element as `True`) and always returns a `bool` tensor. The PyTorch frontend did not produce that `bool` result. The ExportedProgram frontend lowered `logical_and.default` with `self._binary_op(relax.op.logical_and, operator.and_)`, which kept the operand dtype and emitted `relax.op.logical_and` on non-bool inputs (for example `float32`). `relax.op.logical_and` requires boolean inputs and otherwise fails `LegalizeOps` in the TOPI `logical_and`. The FX frontend did not register `logical_and` at all, so the op was unsupported there. ### Changes - Add a shared `_logical_and` converter in `BaseFXGraphImporter` that casts non-bool operands to `bool` before applying `relax.op.logical_and`. Bool operands are passed through unchanged (no redundant cast). - Point the `logical_and.default` (ExportedProgram) registration at the new converter, and add a `logical_and` (FX) registration that was previously missing, matching the existing `logical_not` converter. - Add a standalone `test_logical_and` to both the FX and ExportedProgram test suites asserting the corrected IR (`astype` to bool on each operand, then `logical_and`, producing a `bool` output). ### Notes The cast to `bool` lowers to an elementwise nonzero test, so it matches PyTorch's "nonzero is True" semantics for float, integer, and NaN inputs. (cherry picked from commit 0e32507)
Remove the PyPI-specific ref guard from the wheel publishing workflow. (cherry picked from commit 7151991)
Bump tvmjs version to 0.25.0-dev1 (cherry picked from commit 15b1d98)
This PR try to fix apache#19676 , allow ``CommReduce`` handle 0-dim data correctly. --------- Co-authored-by: flashmouse <flashmosue2012@gmail.com> (cherry picked from commit e067869)
Convert SHA-pinned third-party actions to their full version tags across all workflows and composite actions; each tag verified to resolve to the same commit that was pinned. Drop the unused actions:read / contents:read permissions from the upload_pypi job (aligns with apache/tvm-ffi). No behavior change. (cherry picked from commit 47ddef8)
…che#19706) Two test files under `tests/python/s_tir/meta_schedule/` use the TIRx script dialect (`from tvm.script import tirx as T`) but still call the old TIR block API, which does not exist in `tvm.tirx.script. This pr fixes this (cherry picked from commit fda3220)
This commit fixes and updates the releaase verification script, including guarded binary scan and runtime deps installation. (cherry picked from commit 9b5cd5d)
… favor of default representation (apache#19709) Python-side meta-schedule classes (`PyCostModel`, `PyFeatureExtractor`, `PyMeasureCallback`, `PyScheduleRule`, `PyMutator`, `PyPostproc`) carried an `f_as_string` callback whose only purpose was to produce a repr-style string (`s_tir.meta_schedule.<SubclassName>(0x...)`) for `str(...)`. This mechanism stopped working after apache#19461 migrated `ReprPrinter` to the tvm-ffi `__ffi_repr__` mechanism and intentionally removed the per-type `set_dispatch<Py*Node>` hooks that called back into `f_as_string`, which broke three `*_as_string` tests: - `test_meta_schedule_cost_model.py::test_meta_schedule_cost_model_as_string` - `test_meta_schedule_feature_extractor.py::test_meta_schedule_feature_extractor_as_string` - `test_meta_schedule_measure_callback.py::test_meta_schedule_measure_callback_as_string` Rather than restoring the old behavior, this PR removes the mechanism entirely: the string it produced is just a repr, and tvm-ffi reflection already provides an auto-generated default repr for every object. Keeping a dedicated Python → FFI → C++ callback chain alive only to reproduce that is not worth the complexity. (cherry picked from commit 219f1d8)
This fixes CoreML Relax partitioning after `FoldDataflowBlockOutput` was folded into `CanonicalizeBindings`. `partition_for_coreml` still called `relax.transform.FoldDataflowBlockOutput()`, but that pass is no longer exported. This caused CoreML partition tests to fail with: ```text AttributeError: module 'tvm.relax.transform' has no attribute 'FoldDataflowBlockOutput' ``` The old pass behavior is now covered by `CanonicalizeBindings`, so this updates CoreML partitioning to call `CanonicalizeBindings()` instead. The existing CoreML test file now also includes a partition-only regression test. The end-to-end CoreML tests remain guarded by the CoreML runtime requirements, while the partition test can run without coremltools, Xcode, or a CoreML runtime. (cherry picked from commit 1f7c99d)
…che#19704) Validating the apache-tvm wheel in a minimal environment (no torch, scipy, cloudpickle, or tornado installed) produced 33 pytest collection errors from module-level imports of optional packages. Add pytest.importorskip guards so these modules are reported as skipped instead of erroring during collection. Indirect import chains guarded: - tvm.topi.testing imports scipy - tvm.s_tir.meta_schedule.testing.local_rpc (tvm.rpc.tracker) requires tornado - tvm.s_tir.dlight.benchmark imports cloudpickle Also remove a stray pre-license-header `import pytest` in test_runtime_builtin_paged_attention_kv_cache_flashinfer.py. (cherry picked from commit def37e3)
This fixes CI script tests that accidentally run helper scripts with the
wrong Python interpreter.
Some tests call `run_script(..., env={...})` to pass test-specific
environment variables. `subprocess.run` treats `env` as a full
environment replacement, so the subprocess lost the current `PATH`.
Since these scripts use a shebang like:
```python
#!/usr/bin/env python3
```
they could fall back to the system Python instead of the active test
Python.
In the wheel test environment this caused the CI helper scripts to run
under Python 3.8, even though TVM requires Python >=3.10. Importing
`ci/scripts/jenkins/git_utils.py` then failed on PEP 604 annotations
such as:
```python
Any | None
tuple[str, str] | None
```
with:
```text
TypeError: unsupported operand type(s) for |: '_SpecialForm' and 'NoneType'
```
This patch updates `tests/python/ci/test_utils.py::run_script` to merge
test-provided environment variables into `os.environ` instead of
replacing the entire environment. This preserves the active Python/PATH
while still allowing each test to override or add variables.
(cherry picked from commit 795cd2d)
Newer LLVM versions (observed with LLVM 20) print a scalable broadcast store as a splat constant, e.g. `store <vscale x 4 x float> splat (float 1.000000e+00)`, instead of the older `shufflevector (<vscale x 4 x float> insertelement (...` form. Accept either representation in test_scalable_broadcast so the test passes across LLVM versions while still verifying scalable vector codegen. --------- Co-authored-by: tqchen <tianqi.tchen@gmail.com> (cherry picked from commit e214468)
…apache#19720) Clarify in the Security Guide that loading serialized artifacts must only be done with input from a trusted source. (cherry picked from commit 67627ed)
This PR slims `tvm.libinfo` into a thin *info* layer that delegates path discovery to the `tvm_ffi.libinfo` primitives and never loads libraries. Loading responsibilities move to `tvm.base`, and the various ad-hoc path-finding helpers are phased out in favor of the tvm-ffi resolvers. ## Changes - **libinfo**: add `find_libtvm_runtime()` (resolves `libtvm_runtime` via `_find_library_by_basename` + `_resolve_and_validate`) and `find_tvm_include_path()` (TVM's own `include/`). `find_include_path()` now returns `[find_tvm_include_path(), *tvm_ffi.libinfo.include_paths()]`, folding in the FFI + dlpack + python-helper include dirs. Remove `find_lib_path`, `get_dll_directories`, `use_runtime_lib`, `split_env_var`, and `load_backend_libs`. - **base**: receive `load_backend_libs` and the backend DSO list; the runtime-only switch becomes a strict `TVM_USE_RUNTIME_LIB == "1"` check. - **rpc**: `with_minrpc` uses `find_libtvm_runtime()` (the `runtime` kwarg is retained as an inert back-compat parameter); the rpc server `load_library` resolves the literal library name against the current working directory. - **wasm**: move the `web/dist` asset search into `emcc.find_wasm_lib`, used by `emcc.create_tvmjs_wasm` and the tvmjs asset lookup. - **hexagon**: fix a latent bug where `_get_hexagon_rpc_lib_dir` called a non-existent `tvm_ffi.libinfo.find_lib_path`; it now relies solely on the `HEXAGON_RPC_LIB_DIR` environment variable. (cherry picked from commit ebc9348)
…he#19717) The generic-target tests in test_target_codegen_vulkan.py are auto-parametrized over all enabled targets, which may include nvptx. For nvptx, TVM produces PTX codegen output but not a directly launchable runtime module, so executing the compiled function fails with errors such as cuModuleGetFunction CUDA_ERROR_NOT_FOUND. Add a small helper that skips runtime execution for nvptx after a successful compile, so codegen is still exercised while the invalid runtime launch is avoided. Vulkan-only tests are unchanged. (cherry picked from commit 1e096d6)
This pr fixes two contrib test failures that show up when running source-tree tests against the `apache-tvm` wheel. - Launch the pickle memoize helper script with `sys.executable` so subprocesses use the same Python environment as pytest. - Fix `tvm.contrib.random.random_fill` for packed sub-byte dtypes by using the actual tensor storage size instead of the logical element count. ## Root Cause `tests/python/contrib/test_memoize.py` executed `pickle_memoize_script.py` directly. In a wheel test environment, the shebang can resolve to a different Python than the active wheel venv, causing `ModuleNotFoundError: No module named 'tvm'`. `random_fill` used the product of tensor shape as the number of values to write. For packed dtypes such as `int4`, two logical elements share one byte, so this wrote past the allocated storage and caused native heap corruption / process abort. --------- Co-authored-by: tqchen <tianqi.tchen@gmail.com> (cherry picked from commit 0ff0f85)
…19724) The disco CCL tests called tvm.get_global_func("runtime.disco.compiled_ccl") at module import time, which raises ValueError on TVM builds without Disco CCL support (e.g. published wheels), failing pytest collection. Resolve the function with allow_missing=True and skip the module when it is absent. (cherry picked from commit 3d95dec)
This pr gates the Relax multi-GPU VM test on the actual number of visible CUDA devices. (cherry picked from commit a1d46dc)
…rors (apache#19722) Replace TVM's `Diagnostic` / `DiagnosticContext` machinery with the tvm-ffi `visit_error_context` mechanism. Validators throw an `ffi::Error` seeded with the offending node; leaf pass executors (`ModulePass` / relax `Function` / `DataflowBlock`) catch and rethrow `EnrichPassErrorWithContext`, which appends the failing pass name and a TVMScript-rendered, underlined source location. `relax.analysis.well_formed` now throws on the first violation; a new `check_well_formed` returns a bool, and all C++/Python/test callers are routed accordingly. `include/tvm/ir/diagnostic.h` and `src/ir/diagnostic.cc` are deleted. The enrichment renders with `num_context_lines=10` so a small function shows in full with no skipped-lines marker, while a large module stays bounded. The TVMScript parser diagnostics (`python/tvm/script/parser/core/diagnostics.py`) stay self-contained pure-Python with no `DiagnosticContext` dependency, and restore multi-line source rendering: a diagnostic whose offending AST node spans multiple source lines now renders every spanned line with its gutter line number and an underline covering the span. `tvm.error.DiagnosticError` (used by the TVMScript parser) is retained. A rendered end-to-end enriched-error example is posted as a comment below. (cherry picked from commit 4d28424)
This pr makes the Hexagon pytest plugin avoid importing Hexagon build/session machinery at module import time. This keeps pytest collection lightweight and avoids collection-time failures from Hexagon-specific runtime/session imports when running source-tree tests against a TVM wheel or an environment that does not need to execute Hexagon tests. (cherry picked from commit e41454d)
This pr refactors `pyproject.toml` and deletes redundant dependencies (cherry picked from commit a4a2bf5)
…ble (apache#19730) The remote() helper in tests/python/nightly/test_nnapi/conftest.py returned None when TVM_TRACKER_HOST, TVM_TRACKER_PORT, or RPC_DEVICE_KEY was not set, causing every caller doing `remote_obj, tracker = remote()` to fail with `TypeError: cannot unpack non-iterable NoneType object`. Call pytest.skip() with the list of missing environment variables instead, so the 30 tests in test_ops.py (and test_network.py) report as skipped rather than failed when no NNAPI remote device is configured. Behavior with a configured remote is unchanged. (cherry picked from commit f0ed799)
… let binds to T.let (apache#19729) Fix the s_tir tests broken or left stale by two upstream changes. * test_meta_schedule_space_cuda.py (cap, dil, gmm, t2d, nrm, sfm, cbr, tbg) and test_meta_schedule_space_cuda_async.py (c2d): apache#18927 expanded DefaultCUDA unroll_max_steps from {0, 16, 64, 512, 1024} to {0, 16, 32, 64, 128, 256, 512, 1024} without updating the recorded SampleCategorical decisions. Remap the indices (2->3, 3->6, 4->7) so each test keeps sampling the same unroll value; every sketch was re-verified by replaying the trace and structurally comparing against the expected module. * T.let migration: since apache#19581 the TIRx parser treats `v: T.int32 = expr` as a mutable local-scalar buffer instead of an immutable bind, which is now spelled `v: T.let[T.int32] = expr` (a Bind node, the same form te.create_prim_func emits). Tests whose intent is a bind are migrated to the new spelling: reduction combiner temporaries (add_rfactor, lower_cross_thread_reduction) and let-dependent passes (compact_buffer_region, hoist_expression, remove_undef). * Also convert reduction temporaries in still-green tests (cross_thread_reduction rule, compute_inline, schedule utilities, parallel_vectorize_unroll postproc, dlight general reduction, relax cuda_graph) so the hand-written workloads match the canonical Bind form instead of feeding rules a mutable-scalar body. (cherry picked from commit 67b0c6c)
The all-platform-minimal NDArray runtime test is now superseded by tvm-ffi's own NDArray coverage. This removes the now-redundant `test_runtime_ndarray.py` test file. (cherry picked from commit a2b10a3)
…sage (apache#19748) apache#19714 intentionally removed packed sub-byte (bits 1/4) support from RandomEngine after the int4 fill path was found to write past the allocated storage, but left "int4" in test_random_fill's dtype list, so the test now fails with InternalError on every run. Remove it from the list to match the supported set. Also cast dtype.code / dtype.bits to int in the unsupported-dtype error messages: both fields are uint8_t and were streamed as (invisible) characters, e.g. "Doesn't support dtype code dtype bits". (cherry picked from commit 59f8283)
llvm.experimental.stepvector was renamed to llvm.stepvector in LLVM 20, not LLVM 18. On LLVM 18 and 19 the test picked the new name, which those versions do not know, failing call_llvm_intrin with "Unknown llvm intrinsic function llvm.stepvector". Use the new name only for LLVM >= 20. (cherry picked from commit 2c46ff9)
) This pr marks it xfail with a TODO so the s_tir/transform CI enrollment (apache#19737) is not blocked; the mark should be removed once the CSE determinism fix land. (cherry picked from commit 6e6627e)
…9737) The python-unittest stage enumerates test directories explicitly, and tests/python/s_tir/transform was never enrolled when the transform tests were migrated into the s_tir namespace (apache#18722). As a result the directory has not been exercised by CI at all, which let several regressions land on main unnoticed Add the directory to TEST_FILES so the s_tir transform passes are covered. GPU-marked tests are filtered by the -m gpu marker selection in the gpuonly variant, so CPU enrollment is safe. (cherry picked from commit 6f9609f)
This pr follows the cibuildwheel 4.0 announcement, the official prebuilt
quay.io/manylinux_cuda images ship the CUDA toolkit preinstalled under
/usr/local/cuda on the same manylinux_2_28 baseline the wheel targets.
also switchs the Linux build_cuda_runtime sidecar legs to
quay.io/manylinux_cuda/manylinux_2_28_{x86_64,aarch64}_cuda13_1 and drop
the per-run NVIDIA local-repo RPM download plus dnf install of the
toolkit, which pulled several GB from developer.download.nvidia.com on
every run and was the main source of sidecar build-time variance. This
moves the Linux sidecar from CUDA 13.0.2 to 13.1; the soname stays
libcudart.so.13, so the repair-step glob excludes (libcudart.so.*) are
unaffected and the sidecar carries no device code. The Windows sidecar
keeps using conda and is unchanged.
(cherry picked from commit 0ec3658)
LLVM 20 added vector support for llvm.lround/llvm.llround. In LLVM release/19.x, the IR Verifier still rejects vector operands/results for these intrinsics with "Intrinsic does not support vectors". In release/20.x, that verifier check was removed, and LLVM 20.1 compiles the vectorized lround module successfully. This updates the test gate so LLVM <= 19 still exercises the expected failure path, while LLVM 20+ uses the successful compile path. Co-authored-by: tlopexh <tlopexh@users.noreply.github.com> (cherry picked from commit 235387d)
This fixes the legacy predicated `ptx.cp_async` codegen path used by `InjectPTXAsyncCopy` for `if_then_else(..., 0)` stores. The old inline CUDA emission zero-filled the shared-memory destination when the predicate was false. The TIRx helper-based legacy codegen only skipped the `cp.async`, leaving the destination slot unchanged. This restores the previous behavior by emitting an `@!p st.shared.*` zero store in the generated legacy predicated helper. The CUDA source snapshot in `test_s_tir_transform_inject_ptx_async_copy.py` is updated to reflect the restored false-predicate zero-fill instruction and the current generated helper-based CUDA source. (cherry picked from commit 126f1ba)
…ons (apache#19752) `tests/python/codegen/test_target_codegen_aarch64.py` cross-compiles AArch64 SVE kernels and regex-matches the generated assembly for specific instruction forms and counts. Several of those assertions encode the exact code shape produced by the LLVM versions used in CI (15-17). On a TVM built against LLVM 20 the tests fail, even though the emitted IR is correct (+sve target-features and vscale_range(1,16) are present) -- the difference is entirely inside LLVM's loop vectorizer / cost model, not in TVM's codegen. (cherry picked from commit 383c446)
) ### Motivation `torch.logical_or` and `torch.logical_xor` accept input tensors of any dtype (treating any nonzero element as `True`) and always return a `bool` tensor. Neither op was handled by the PyTorch frontend. The ExportedProgram frontend did not register `logical_or.default` / `logical_xor.default`, and the FX frontend did not register `logical_or` / `logical_xor`, so importing a model that uses either op failed early with `Unsupported function types`. This follows up on apache#19679 (`logical_and`) and addresses the explicit question raised in apache#19743: whether `logical_or` and `logical_xor` need the same handling. ### Changes - Add shared `_logical_or` and `_logical_xor` converters in `BaseFXGraphImporter` that cast non-bool operands to `bool` before applying `relax.op.logical_or` / `relax.op.logical_xor`. Bool operands are passed through unchanged (no redundant cast). - Register `logical_or.default` / `logical_xor.default` (ExportedProgram) and `logical_or` / `logical_xor` (FX), matching the existing `logical_and` converter. - Add standalone `test_logical_or` and `test_logical_xor` to both the FX and ExportedProgram test suites, asserting the corrected IR (`astype` to bool on each operand, then the logical op, producing a `bool` output). ### Notes The cast to `bool` lowers to an elementwise nonzero test, so it matches PyTorch's "nonzero is True" semantics for float, integer, and NaN inputs. (cherry picked from commit f5e8a3a)
…s, gemm-async, backend reorg (apache#19757) This PR batches several post-bringup TIRx follow-ups, rebased onto current `main`. ### Changes - **op-dispatch**: per-call exec scope via `Tx.<scope>.op`; remove `ExecScopeStmt` - **namespaces**: split TIRx op namespaces; remove tile-primitive kind attrs - **codegen**: support explicit CUDA launch bounds - **gemm-async**: support contiguous-axis (K-major) operand slicing - **backend reorg**: move in-tree GPU backends out of core into `src/backend/<target>/` and `python/tvm/backend/<target>/` (codegen/runtime/op), with the corresponding `CMakeLists.txt` / `cmake/modules` and include-path updates ### Testing - Builds with `USE_CUDA=ON` / `USE_LLVM=ON` - The TIRx Python test suite (`tests/python/tirx/`) passes locally (cherry picked from commit bb6f8ae)
…he#19758) CUDA graph VM support is Relax VM functionality, so it should stay with the VM runtime layout instead of living under the CUDA backend runtime tree introduced by the backend reorg. This moves the CUDA graph builtin back under src/runtime/vm/cuda and points the CUDA runtime CMake source glob at that VM-owned location while preserving the existing symbols and behavior. (cherry picked from commit d291513)
The CoreML runtime still used legacy packed argument conversions, including direct AnyView-to-string conversion, type_code checks, and old tensor handle constants. After the FFI update, enabling USE_COREML failed to compile and the runtime create function was not registered. This pr updates the runtime to use current PackedArgs cast helpers for strings, indices, and DLTensor arguments, and use the Objective-C NSArray type for CoreML metadata inputs. This restores the CoreML runtime registration when USE_COREML is enabled. (cherry picked from commit 07fe3fe)
…ecar (apache#19761) This pr is the follow-up pr of apache#19754. It removes the redundant `pip install -U pip cmake ninja` (since those are already installed in docker) and the unused python_bin variable, and update the comment to match: the image provides the build tools, and PATH only needs the bundled CPython (for the configure) plus the CUDA toolchain (for nvcc). Also fixs a stale comment in the Windows sidecar: it claimed the conda 13.0.2 pin matched the Linux hook, but the Linux side is now CUDA 13.1 (from the image). The win-64 nvidia channel still caps at 13.0.x, so Windows stays on 13.0.2 -- slightly behind Linux but harmless, since the sidecar has no device code and links the CUDA runtime by soname only. The Windows pip install is kept: that build uses the Ninja generator and runs on a runner without the prebuilt CUDA image. (cherry picked from commit 1c27ccf)
…#19760) ## Summary The in-tree target custom datatype path adds maintenance surface while current development focuses on core datatypes. This PR phases out the built-in registry/lowering implementation and keeps the core dtype behavior intact. - Remove the target/datatype implementation, BYODT posit build option, and related Python helpers - Remove the custom datatype lowering pass from TIRX and S-TIR finalization pipelines - Simplify remaining TIRX dtype handling back to built-in/core datatypes (cherry picked from commit e43555f)
The backend reorg leaves target-kind and op registration behind generic register.cc shims, which makes the new layout harder to navigate and keeps a few obsolete target paths alive. This change keeps backend registration beside the backend definitions while preserving the existing static initialization behavior. Backend target-kind registration units now use semantic target_kind.cc names, target builtin registration is colocated with the CUDA, Metal, and Trainium builtin definitions, the NVPTX intrinsic rules live with CUDA LLVM codegen, and the unused target/opt placeholder is removed. (cherry picked from commit 8b79dfd)
…#19767) Adds on-screen info about samples failed to build, that may happen in early stages. ---------- #### Before: ``` 2026-06-14 18:36:30 [INFO] [task_scheduler.cc:196] TaskScheduler picks Task #0: "dense" 2026-06-14 18:36:32 [INFO] [task_scheduler.cc:209] Sending 64 sample(s) to builder 2026-06-14 18:38:47 [INFO] [task_scheduler.cc:211] Sending 64 sample(s) to runner ``` #### After: ``` 2026-06-14 19:53:23 [INFO] [task_scheduler.cc:193] TaskScheduler picks Task #0: "dense" 2026-06-14 19:53:25 [INFO] [task_scheduler.cc:206] Sending 64 sample(s) to builder 2026-06-14 19:55:39 [INFO] [task_scheduler.cc:215] Build errors: 64 sample(s) 2026-06-14 19:55:39 [INFO] [task_scheduler.cc:217] Sending 0 valid sample(s) to runner ``` (cherry picked from commit 4b7b7de)
The CUDA-specific timer preprocessor test was the only user of the nvbench-derived L2 cache flush helper. Removing that path avoids carrying a special CUDA registration and third-party header for benchmark-only test coverage. This PR removes the CUDA L2 cache flush registration, the dedicated preprocessor timing test, and the now-unused nvbench-derived license/dependency entry. (cherry picked from commit 4119c9f)
…ype (apache#19772) ### Root cause In the ONNX `LayerNormalization` spec the bias `B` is optional; when omitted it should behave as zeros shaped and typed like the scale `W`. In `LayerNormalization._impl_v17`, the synthesized zero bias instead took its shape from `data.struct_info.shape[1]` (an unrelated data dim) and hardcoded `dtype="float32"`. For input `[2, 3, 4, 8]` with scale `[8]` and `axis=-1` this builds a bias of shape `(3,)` while gamma is `(8,)`, so `relax.op.nn.layer_norm` raises a size-mismatch `InternalError`. The float32 hardcode also breaks fp16/bf16 no-bias models, since gamma, beta, and data must share a dtype. PyTorch's `nn.LayerNorm(..., bias=False)` exports exactly this no-bias form. ### Fix Derive both the shape and dtype of the synthesized zero bias from the scale, matching the ONNX semantics for an omitted `B` and the existing torch frontend (`relax.const(np.zeros(shape), x.struct_info.dtype)`): ```python if bias is None: bias = relax.const(_np.zeros(gamma_shape, dtype=scale.struct_info.dtype)) ``` `gamma_shape` and the `_np`/`get_const_tuple` imports are already present. Deriving the dtype from the scale (rather than the issue's float32-only suggestion) is what also fixes the fp16/bf16 case. ### Test plan Added non-square no-bias regression cases to `test_frontend_onnx.py::test_layer_norm` (the previous no-bias case was square, which masked the bug): float32 `[2,3,4,8]`/scale `[8]` and float16 with full `check_correctness`, plus a bf16 importer-only case (ORT's CPU provider has no bf16 LayerNormalization kernel). Fixes apache#19691 (cherry picked from commit 44b55a0)
### Root Cause
The ONNX `CumSum` converter in the Relax frontend rejected any node with
`exclusive=1` via a bare
`assert not attr.get("exclusive", False), "Exclusive option not yet
supported."`, so importing a
model that used exclusive cumulative sums failed with `AssertionError:
Exclusive option not yet
supported.`. The underlying op already supports the exclusive form:
`relax.op.cumsum` forwards an
`exclusive` flag to the FFI and `topi/scan.py` implements the exclusive
branch, so the converter
only needed to pass the attribute through.
### Solution
Drop the assert and read the attribute as `exclusive =
attr.get("exclusive", 0) != 0` (matching the
existing `attr.get("reverse", 0) != 0` idiom in the same converter),
then pass it to
`relax.op.cumsum(data, axis, exclusive=exclusive)`. The existing reverse
handling
(`flip -> cumsum -> flip`) composes correctly with exclusive, so
`reverse=1, exclusive=1` lowers to
an exclusive scan over the reversed axis.
### Test Plan
Extended `test_cumsum` in `tests/python/relax/test_frontend_onnx.py` to
parametrize `exclusive` over
`[True, False]`, so `check_correctness` now exercises all four
`(reverse, exclusive)` combinations
against ONNX Runtime:
```
python -m pytest tests/python/relax/test_frontend_onnx.py::test_cumsum -v
```
### Issue
Fixes apache#19692
(cherry picked from commit 21298c7)
This fixes several race conditions and enhances printing infos:
* Protect against race condition if ```list``` or ```remove``` on
nonexistent filepath is attempted
* ```ListDir``` returns gracefully when folder nonexistent without
exception throwed (blocking)
* ~Make sure ```TVM_FFI_THROW``` also ```_exit()```, not blocking the
child (until timeout)~
* Enhance printed info during ```tvm_rpc``` with concise & trackable
status updates
* Print spent time in the rpc session, may help users setting the
timeout to a safer boundary
This enhanches significantly the speed of remote tuning, previously
stalled due to timeout conditions.
---
#### Issue description:
If worker child process timeouts it gets killed (expected, based on
timeout) leaves with cache file leftovers. Thats why when a new session
starts it will do a cleanup first. But when previous child process
finished normally (means also cleaned up itself), the new cleanup
process will hang due to throwed exception (blocking) over the
"nonexistent directory" error !
In this vicious circle we always (almost) will have hanged child workers
that needs to be killed based on a timeout.
---
#### Before
* Child worker never (mostly) exited on its own, it was always killed
(on timeout):
```
[02:52:48] {..}/rpc_server.cc:310: Connection success 192.168.1.2:39460
[02:52:48] {..}/rpc_env.cc:149: Load module from /tmp//.cache/tvm_tmp_mod.tar.so ...
[02:52:50] {..}/rpc_server.cc:208: Child pid=3727459 killed (timeout = 2), Process status = 15
[02:52:50] {..}/rpc_server.cc:239: Socket Connection Closed
```
#### After
* Based on child condition and its exit we have:
Worker child exits on its own (finish normally, no timeout):
```
[11:58:27] {..}/rpc_server.cc:312: New session from 192.168.1.2:56030
[11:58:27] {..}/rpc_env.cc:123: Load module from
/tmp//.cache/tvm_tmp_mod.tar.so ...
[11:58:28] {..}/rpc_server.cc:337: Finished serving 192.168.1.2:56030
after 1.028 sec
[11:58:28] {..}/rpc_server.cc:212: Child pid=616766 finished, status = 0
[11:58:28] {..}/rpc_server.cc:241: End session with 192.168.1.2:56030
```
Worker child is terminated (unfinished task, due to timeout):
```
[11:23:11] {..}/rpc_server.cc:312: New session from 192.168.1.2:58152
[11:23:11] {..}/rpc_env.cc:123: Load module from
/tmp//.cache/tvm_tmp_mod.tar.so ...
[11:29:26] {..}/rpc_server.cc:208: Child pid=414284 killed (timeout = 15
sec), status = 15
[11:29:26] {..}/rpc_server.cc:241: End session with 192.168.1.2:58152
```
(cherry picked from commit 9831966)
## Summary This PR upgrades TVM's core CMake build baseline from C++17 to C++20. It updates the top-level CMake C++/CUDA standards and the LLVM compile probe fallback flags. Downstream app, JVM, web, Python helper, and vendored third-party build surfaces are intentionally left unchanged so they can migrate separately. ## Verification - Configured a minimal local build with GCC 13.3 and NVCC 13.2: `cmake -S . -B build/cxx20-check -GNinja -DBUILD_TESTING=OFF -DUSE_LLVM=OFF -DUSE_CUDA=OFF -DTVM_BUILD_PYTHON_MODULE=OFF` (cherry picked from commit c9a3c2f)
## Summary The CUDA runtime can rely on the shared tvm-ffi CUDA error helper instead of carrying a TVM-local common header. This keeps CUDA error handling aligned with the FFI CUDA support and removes the remaining reason for src/backend/cuda/runtime/cuda_common.h. Main changes: - Move CUDA workspace thread-local state into cuda_device_api.cc. - Replace CUDA_CALL use sites with TVM_FFI_CHECK_CUDA_ERROR from tvm/ffi/extra/cuda/base.h. - Inline the remaining CUDA driver checks at their call sites and delete cuda_common.h. (cherry picked from commit 5550fb3)
6615ea2 to
7ae1ae7
Compare
## Summary Backend loading is easier to maintain when native backend library discovery, in-tree backend Python hook loading, and out-of-tree entry point autoload are owned by the backend namespace. This PR consolidates those paths under `tvm.backend._autoload_backends` while preserving compatibility routes from the previous top-level helper and `tvm.base.load_backend_libs`. - Move backend runtime DSO loading into `tvm.backend._autoload_backends` - Route `backend.load_all()` through the backend autoload helper - Keep the previous top-level `_autoload_backends` module as a thin compatibility import
Brings the
v0.25.0release branch up tomainbycherry-picking, in order, the commits that landed after rc0. Each commit
keeps its original authorship and a
(cherry picked from commit …)line.Motivation: rc0 is missing recent changes we want in the release — notably the
out-of-tree backend autoload (#19727), the tvm-ffi object/JSON migrations and
per-backend runtime DSOs, and a batch of LLVM-version test gates that keep CI
green.
Notes:
branch stays linear.
follow as a separate step.