## Summary
Compiler Targets can carry device-type semantics that runtime
device-name parsing does not preserve.
- add `tvm.device_from_target` for canonical Target-to-Device
translation
- use explicit runtime constructors where the device kind is fixed
- update target-derived utilities, tests, and documentation to use the
explicit boundary
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.
## Rationale
`SizeVar` encodes nonnegativity in runtime subtype identity, which is
fragile under cloning and remapping. Symbolic integer values should use
one `Var` representation, with nonnegative facts recorded in the
analyzer at the use sites that establish them.
## Changes
- Remove `SizeVar` from the C++, Python, TE, TVMScript, FFI, visitor,
and serialization surfaces, and migrate callers to `Var`.
- Preserve the existing Relax constraint ownership model and use
`MarkGlobalNonNegValue` as the canonical path for global nonnegative
facts.
- Preserve `T.handle()` as the normal opaque-handle form. An optional
dtype constructs a typed pointer, with `T.handle("void")` reserved for
an explicit pointer-to-void.
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`
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.
## Summary
This PR adds the initial TIRx support needed for low-level programming
of Blackwell-class GPU architectures. As part of the ongoing TIRx
refactor, it introduces TVMScript support for directly scripting
advanced hardware features without relying on scheduling as the primary
programming interface.
The change keeps existing `s_tir` script support intact while making
direct scripting a first-class path for TIRx programs.
## Main Changes
- Add TIRx operator dispatch and layout infrastructure.
- Add TVMScript support for new low-level TIRx operations.
- Add analysis, transform, and lowering support for TIRx IR nodes.
- Add CUDA/Blackwell-oriented codegen and intrinsic coverage.
- Add Python and C++ integration points for TIRx scripting and runtime
support.
## Validation
- `pre-commit run --all-files`
- `ninja -C build -j32`
- `CUDA_VISIBLE_DEVICES=2 pytest tests/python/tirx/ -n 16`
- `1723 passed, 47 skipped, 32 warnings`
- `CUDA_VISIBLE_DEVICES=2 python -m pytest -v
tests/python/all-platform-minimal-test`
- `37 passed, 105 skipped`
- `TVM_TEST_TARGETS=llvm python -m pytest -v tests/python/tirx-analysis
tests/python/tirx-base tests/python/tirx-transform -n 16`
- `664 passed, 25 skipped, 9 xfailed, 1 xpassed`
## Local CI Notes
Some full CI-equivalent jobs were not locally reproducible because this
machine is missing parts of the Apache TVM CI environment, including
`llvm-config-15/17`, Vulkan, ROCm, Maven, Sphinx, Doxygen, Emscripten,
and ARM/QEMU cross-toolchain components. Metal-specific tests were
skipped locally because no Metal runtime is available.
## Why
This refactor reshapes each backend into a self-contained
`src/target/<X>/`
cluster (with optional `src/target/<X>/llvm/` for LLVM-dependent
codegen) and
introduces a per-backend fallback module that absorbs the cross-compile
role
cleanly — without `target/opt/` stubs, without `DeviceSourceModuleNode`,
and
without leaking a synthetic `kind()` to consumers.
## High-level principles
- **One directory per backend.** All codegen-side files for backend
`<X>` live
under `src/target/<X>/`. Optional `src/target/<X>/llvm/` subdir for
files
that require `USE_LLVM` at build time. Backend grouping wins over
build-dependency grouping (the latter being upstream's `target/llvm/`).
- **Plugin-only runtime modules.** `src/runtime/<X>/<X>_module.h` is
deleted.
The runtime's real `<X>ModuleNode` is reachable only via the FFI
registry
(`ffi.Module.create.<kind>`, `ffi.Module.load_from_bytes.<kind>`). No
C++ API surface other than the static registrations.
- **Per-backend fallback module for cross-compile.** Each `<X>` gets a
`<X>FallbackModuleNode` in `src/target/<X>/<X>_fallback_module.{h,cc}`.
Same `kind()` as the real module. Codegen-time only — never reachable
via
load. `GetFunction` errors with a backend-specific "runtime not linked"
message; `InspectSource` works.
- **Codegen-side wrapper does the fallback selection.** Codegen calls
`<X>ModuleCreateWithFallback(...)`, which tries
`ffi.Module.create.<kind>`
via the registry; on miss, falls through to `<X>FallbackModuleCreate`
(plain C++; reachable directly from the fallback header). When
`USE_<X>=ON`
is in effect, the registry hit returns the real module; when
`USE_<X>=OFF`, the fallback is what codegen gets. No CMake `if/else`
gating; fallback always compiled.
## Specific changes
### New per-backend directories (codegen + fallback)
- `src/target/cuda/` — `codegen_cuda.cc` + `intrin_rule_cuda.cc` +
fallback module pair + `llvm/codegen_nvptx.cc`
- `src/target/rocm/` — fallback module pair + `llvm/codegen_amdgpu.cc` +
`llvm/intrin_rule_rocm.cc`
- `src/target/hexagon/` — fallback module pair +
`llvm/codegen_hexagon.cc` + `llvm/intrin_rule_hexagon.cc`
- `src/target/metal/` — `codegen_metal.cc` + `intrin_rule_metal.cc` +
fallback module pair
- `src/target/vulkan/` — `build_vulkan.cc` + the rest of `target/spirv/`
absorbed + fallback module pair
- `src/target/opencl/` — `codegen_opencl.cc` + `intrin_rule_opencl.cc` +
fallback module pair
- `src/target/webgpu/` — `codegen_webgpu.cc` + fallback module pair (was
`WebGPUSourceModuleNode`, renamed)
### New fallback classes
`<X>FallbackModuleNode` for X in {`CUDA`, `ROCm`, `Hexagon`, `Metal`,
`Vulkan`, `OpenCL`, `WebGPU`}. Each:
- `kind()` matches the real backend
- Stores `(code or smap, fmt, fmap, source)` — no driver/runtime calls
- `GetFunction` errors with backend-specific "runtime not linked"
message
- `InspectSource` works fully
- `SaveToBytes` byte-identical to real
SPIR-V codegen currently emits `ArrayStride` and `Offset` decorations
for non-interface allocations in `GetStructArrayType()`. That is correct
for descriptor-backed interface blocks, but not for static workgroup
allocations.
I hit this while bringing up tilelang vulkan shared memory allocation
path: the vulkan validation rejected shaders that used shared memory
lowered through this path:
```
tvm.error.InternalError: Check failed: res == SPV_SUCCESS (-10 vs. 0) :
index=44 error:[VUID-StandaloneSpirv-None-10684] Invalid explicit layout decorations on type for operand '25[%_ptr_Workgroup__struct_24]'
%A_shared = OpVariable %_ptr_Workgroup__struct_24 Workgroup
```
FIX: This PR keeps layout decoration for interface blocks, and skips for
non-interface allocations such as static shared/workgroup memory. A new
compile-only test is added for this.
One possible concern is that there's already a pre-existing test using
`fetch_to_shared`.
This PR brings up the tirx namespace. We have been spliting out the
original tir namespace to include high-level component s_tir and this PR
updates the remaining low-level part as tirx namespace
## Summary
- Remove `body` field from `AllocBufferNode` and `DeclBufferNode`,
making them flat statements consistent with `Bind`
- Buffer scope extends to end of enclosing scope via flat `SeqStmt`
semantics
- 60 files changed across core IR, codegen backends, transforms, script
IR builder, and tests
## Test plan
- All existing test suites pass (tir-transform, tir-base, tvmscript,
s_tir, codegen, C++)
## Summary
This PR introduces `AllocBufferNode`/`AllocBuffer` as a single TIR
statement that both allocates memory and declares a buffer into scope.
This replaces the previous pattern of `Allocate(var, dtype, shape, cond,
DeclBuffer(buf, body))` with the simpler `AllocBuffer(buf, body)`.
### Main changes
- **New IR node** `AllocBufferNode` with fields `{buffer, annotations,
body}` — same semantics as `DeclBuffer` but also allocates memory
- **TVMScript**: `T.alloc_buffer(shape, dtype, scope)` now emits
`AllocBuffer` directly (statement-level allocation).
`T.sblock_alloc_buffer(...)` for SBlock-level buffer allocation (full
parameter set)
- **All codegen backends** (C, CUDA, Metal, OpenCL, WebGPU, LLVM, NVPTX,
AMDGPU, SPIR-V) updated to handle `AllocBufferNode`
- **All TIR transforms** (storage_rewrite, flatten_buffer,
vectorize_loop, lower_warp_memory, etc.) updated
- **All S-TIR transforms** (compact_buffer_region, merge_shared_memory,
inject_double_buffer, etc.) updated
- **Removed `AllocateNode`** entirely — `AllocBuffer` is now the sole
allocation primitive
- **Removed `AllocDescriptor`** from merge_shared_memory_allocations —
uses `Buffer` objects directly
- **Added `AllocBuffer::ConstantAllocationSize()`** inline helper method
### Design rationale
The old `Allocate + DeclBuffer` pair was a historical artifact:
`AllocateNode` stored raw fields (`buffer_var`, `dtype`, `extents`,
`condition`) separate from the `Buffer` object, requiring pattern
matching (`IsAllocateDeclBufferPattern`) to reconstruct the buffer
association. `AllocBuffer` unifies this into a single node with a proper
`Buffer` reference, simplifying codegen backends and transform passes.
225 files changed, ~3500 insertions/deletions (net near-zero, mostly
mechanical migration).
## Test plan
- [x] All TIR base tests pass
- [x] All TIR transform tests pass
- [x] TVMScript roundtrip tests pass
- [x] S-TIR transform tests pass
- [x] Codegen tests pass
- [x] All-platform minimal tests pass
- [x] C++ functor tests pass
- [x] Pre-commit clean (clang-format, ruff, etc.)
This PR phases out legacy target string format in favor of the json
style format that is more well formed. It also simplfies our overall
code in handling multiple formats.
This PR migrates all the codegen tests to explicitly using tvmscript
instead of indirectly via s_tir.Schedule. They makes the test surface
more unit, contains less dep and more maintainable.
This PR initalizes the s_tir for scheduable TensorIR. The change mainly
starts from python side, the we will gradually move towards the c++ side
in followup PRs. The python main change:
tir.Schedule => s_tir.Schedule
This PR renames tir.Block to SBlock. This clearly indicate the
scheduable property of the block and is a prereq for followup stir
passes refactor.
Main changes:
- Data structure change from Block to SBlock
- Syntax change from T.block to T.sblock
This PR Updates the NDArray => Tensor.
Both tensor and ndarray are commonly used terms.
Because the term Tensor is getting more common in the context of ML,
we do the rename to stay more aligned with torch.Tensor and DLTensor.
This PR formalizes original runtime::Module into ffi
as ffi.Module and cleans the APIs around it.
The goal is to stablize the Module API as extra API that can benefit the overall
ffi interactions. We also refactors the c++ code that depends on the Module.
* [REFACTOR] Phase out te.schedule python components
This PR phases out te.schedule python components.
te.compute is kept around for future usages.
tir.Schedule is a more modern version of the scheduling that we can use onwards.
Doing so also helps us to cleanup the testcases that relies on
explicit full build and execution. As we move future unit testcases
towards structural equality based unit tests.
* Simplify CI to focus on UT
The main rationale is that we should only have very few target
dependent UT in tests/python/codegen and possible
a new category in future for op-level integration if needed.
* Re-enable wasm
* fix lint
* remove hybrid,sparse autodoc and remove tests
---------
Co-authored-by: Siyuan Feng <hzfengsy@sjtu.edu.cn>
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>
* use `20241105-030952-3e386fd3` for ci
* fix missing `language` kwarg in save_rst_example
* disable invalid-name check by pylint
* disable caffe frontend tests
* disable mxnet frontend tests
* fix stablehlo importer
* enable some stablehlo tests
* skip tests because mxnet raise AttributeError
* use stable sort for `np.argsort`
* fix `TypeError: arrays to stack must be passed as a "sequence" type such as list or tuple.`
* disable tests because of shape error
* remove `get_html_theme_path` because it's deprecated
* ignore warnings from sphinx
* disable oneflow frontend tests
* remove oneflow tutorial
* remove debug print
The current unittest folder is too large and contains too many files and
too many components. This PR refactors the unittest folder by moving the
files to the corresponding folders.
Tianqi Chen