Fixes#19887.
This PR fixes several Relax TensorRT BYOC issues exposed by YOLO-style
models:
- adds TensorRT support for SiLU and resize2d
- preserves operand and TupleGetItem ordering during codegen
- fixes cyclic and unsafe Tuple/TGI region merging
- handles static Shape bindings and nested packed-function outputs
- normalizes PrimType dtype arguments passed to relax.arange
With these changes, yolo11n-seg can be merged into a single TensorRT
region, while yolo11n can be imported and partitioned successfully.
## 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.
## Summary
- Make `PrimExpr` a typed C++ view over `Expr` values whose
`ExprNode::ty` is `PrimType`, instead of using a separate runtime node
class as the proof of primitive-ness.
- Use the shared `ir::Call` node for Relax, TIRX, and primitive-valued
calls, while keeping primitive-only APIs explicit at their semantic
boundaries.
- Keep Python on the general `Expr` surface for primitive-typed values
so `isinstance` behavior does not imply a nominal primitive-expression
subclass.
## Design Rationale
The main advantage of this change is that common expression nodes such
as `Call` can be unified without specializing each one to `PrimType`. A
single `ir::Call` can represent a Relax tensor call, a Relax scalar
call, or a primitive-valued intrinsic call; the result type stored in
`ExprNode::ty` determines whether that particular value can be viewed as
`PrimExpr`.
This keeps the IR node hierarchy focused on expression structure rather
than result-type categories. Nodes that are intrinsically primitive,
such as integer and floating-point literals or TIRX primitive operators,
still have strongly typed C++ APIs and data structures. General nodes
whose result type may vary, such as `Call`, remain general `Expr` nodes
and are narrowed to `PrimExpr` only where primitive-only semantics are
required.
The PR also keeps the compatibility surface practical: C++
primitive-only APIs continue to accept `PrimExpr`, Python exposes a
compatibility predicate for checking the primitive typed category, and
visitors/printers use one natural `Call` path rather than duplicating
Relax and primitive call handling. Missing expression types are
represented explicitly with `Type::Missing()` so constructors can leave
type inference to later analysis without relying on nullable `Type`
values.
This PR lets Relax expressions directly take `PrimExpr` values without
requiring the explicit `PrimValue` wrapper, continuing the Relax IR
unification work by removing Relax-specific leaf/base expression layers.
Summary:
- Remove `LeafExpr` / `LeafExprNode` and use direct expression-node
checks where needed.
- Converge Relax expression typing onto the shared IR `Expr` base.
- Remove the `PrimValue` node wrapper while keeping `relax.prim_value` /
`R.prim_value` as conversion helpers that return existing `PrimExpr`
values unchanged.
- Register direct `PrimExpr` handling through exact concrete node
dispatch, aligned with the `tirx` expression visitor list and excluding
arith iter-map intermediate nodes.
- Inline the private Python primitive conversion helper into public
`relax.prim_value`.
- Handle direct `PrimExpr` values in frontend scalar paths without
assuming a `.value` field on non-immediate expressions.
## Summary
- unify Relax's former StructInfo surface into the Type vocabulary and
Expr.ty storage path
- remove leftover DependentTypeNode and legacy OpNode::op_type storage
- keep base Type nullable while concrete Relax/DTensor type refs are
non-nullable
- clean stale StructInfo/TensorStructInfo/sinfo vocabulary in
Python/docs and distributed-op macros
- address Gemini follow-ups for parser annotations, BlockBuilder
docstring, and Adreno TensorType cast audit
This pr is the follow-up pr to #19810 Add partition_for_tensorrt, which
offloads TensorRT-supported subgraphs from a module with a single call,
together with the pattern table whose composite names ("tensorrt.<op>")
match the runtime converter registry. This is the entry point other BYOC
backends expose as partition_for_<name>.
This pr is the follow-up pr to #19789. CurrentTensorRT BYOC converters
were ported from Relay and still read attribute names/shapes that no
longer match the Relax ops, so most ops crashed ("Key: <name> is not
found") or produced wrong results when offloaded.
This pr changed
- Converters (tensorrt_ops.cc): port reduce, matmul, expand_dims,
layer_norm, clip, reshape, strided_slice, split and layout_transform to
read Relax's attributes/arguments. Notable shape changes: clip min/max
are PrimValue arguments (not a_min/a_max attrs), reshape's shape is a
Shape argument, matmul has no transpose flags, split is multi-output
with no "mode", and layout_transform is an IndexMap rather than
src/dst_layout strings. Unsupported cases (non-static reshape,
non-permutation layout_transform) now raise a clear error instead of
crashing.
- Codegen (codegen.cc): serialize an op's non-tensor arguments
(PrimValue / ShapeExpr / tuple) as "arg_"-prefixed node attributes,
materialize a reduce op's all-axes default, and translate a
pure-permutation layout_transform IndexMap into a transpose order.
- Runtime: disable the TF32 builder flag so offloaded FP32 subgraphs
match TVM's FP32 reference, and use a process-lifetime TensorRT logger
(a per-runtime logger was left dangling once its runtime was destroyed,
corrupting the heap during TensorRT teardown).
All tests are validated locally.
This pr fixes#19609. TensorRT 10 removed a large set of APIs that the
Relax TensorRT BYOC integration relied on, so it failed to compile
against TRT >= 10. Port the runtime and codegen to the TRT10 API and
require TensorRT >= 10:
- Lifetime: obj->destroy() -> delete (destroy() removed in TRT10).
- Builder: drop implicit-batch mode (networks are always explicit-batch
via createNetworkV2(0); setMaxBatchSize removed); setMaxWorkspaceSize ->
setMemoryPoolLimit(kWORKSPACE); buildEngineWithConfig ->
buildSerializedNetwork + deserializeCudaEngine, keeping the IRuntime
alive alongside the engine.
- Execution: the binding-index model (getNbBindings / getBindingIndex /
setBindingDimensions / execute / executeV2) -> the named-tensor model
(getNbIOTensors / setInputShape / setTensorAddress / enqueueV3);
deserializeCudaEngine drops the trailing IPluginFactory* argument.
- Layers: addConvolution / addPooling / addDeconvolution / addPadding ->
the *Nd variants; set{Stride,Dilation} -> *Nd; IFullyConnectedLayer /
addFullyConnected removed -> dense rebuilt with addConstant +
addMatrixMultiply.
- Add a build-time guard that emits a clear error on TensorRT < 10.
Also fix pre-existing issues that prevented this path from running
end-to-end: the runtime had drifted from the current tvm-ffi API
(TVMTensorCopyToBytes / TVMGetLastError, VectorToTrtDims over
ffi::Array, a stale `override` on the destructor), and the conv
converters read a Relay-era "channels" attribute that Relax does not
emit (output channels are now derived from the kernel shape).
All tests are verified correct locally. This pr barely includes api
updates and there is no new parts added
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`
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.
Prior to this commit, `relax.transform.LegalizeOps` needed to be
called prior to `relax.build`. This commit adds `LegalizeOps` to the
lowering flow, to simplify the calling steps for an end-user. If the
`IRModule` contains no legalizable functions, a second legalization
pass has no effect.
Some test cases relied on this behavior as an implicit assertion that
operator fusion patterns applied. That is, by omitting `LegalizeOps`,
a successful compilation `relax.build` would only occur if all
legalizable operators have already been removed, and so an incorrect
fusion pattern would result in a failure to build the module. While
these tests would be better expressed by comparing against an expected
fused pattern, updating the tests is outside the scope of this PR. To
allow these tests to keep their implicit assertions, a
`"relax.transform.apply_legalize_ops"` config can be used to disable
the `LegalizeOps` pass.
This PR refactors relax build so it get exposed at the opt-level.
We also introduces an explicit jit functionality to handle
live loading of compiled artifacts from cutlass.
We also move relax vm to runtime so it can be clearly isolated
from the rest of the compiler stack.
Shushi Hong