[TIRx][Runtime] Rebase TVM to 0.26.dev0 for TileLang#50
Open
LeiWang1999 wants to merge 106 commits into
Open
Conversation
…che#19495) This PR adds first-class `nn.ParameterList` and `nn.ParameterDict` containers to the Relax frontend. These containers provide PyTorch-like list/dict registration for raw `nn.Parameter` objects while preserving Relax frontend semantics: values must be explicit `nn.Parameter` instances, with no automatic tensor-to-parameter conversion. ### Changes - Add public `nn.ParameterList` and `nn.ParameterDict` exports. - Support stable parameter names in traversal: - `params.0`, `params.1` - `params.foo`, `params.bar` - Integrate the new containers with: - `named_parameters()` - `parameters()` - `state_dict()` - `load_state_dict()` - `to(dtype=...)` - `export_tvm()` - `nn.Mutator` - Add focused tests for basic container behavior, type validation, nested traversal, export parameter names, state loading, dtype conversion, and mutator naming. --------- Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com> (cherry picked from commit 86794e7)
## Summary This PR adds Relax TFLite frontend support for the following segment operators from apache#19412: - `SEGMENT_SUM` - `UNSORTED_SEGMENT_MIN` - `UNSORTED_SEGMENT_PROD` These operators are lowered through `relax.op.scatter_nd` with the corresponding reduction modes. ## Changes ### TFLite Frontend 1. Add TFLite converter mappings for segment operators: - `SEGMENT_SUM` -> `scatter_nd(..., reduction="add")` - `UNSORTED_SEGMENT_MIN` -> `scatter_nd(..., reduction="min")` - `UNSORTED_SEGMENT_PROD` -> `scatter_nd(..., reduction="mul")` 2. Add shared segment lowering logic: - Convert `segment_ids` into scatter indices via `expand_dims`. - Build the output shape from `num_segments` or constant `segment_ids`. - Initialize the scatter base tensor with the correct reduction identity. ### Tests Add TFLite frontend tests for: - `test_segment_sum` - `test_unsorted_segment_min` - `test_unsorted_segment_prod` Each test verifies the imported Relax IR lowers to `R.scatter_nd` with the expected reduction mode and base tensor initialization. ## Testing All targeted tests pass: ```bash python -m pytest \ tests/python/relax/test_frontend_tflite.py::test_scatter_nd \ tests/python/relax/test_frontend_tflite.py::test_segment_sum \ tests/python/relax/test_frontend_tflite.py::test_unsorted_segment_min \ tests/python/relax/test_frontend_tflite.py::test_unsorted_segment_prod \ -q ``` ## References - Issue apache#19412: TFLite Relax frontend operator support tracking - Related PR apache#19490: Adds SCATTER_ND support (cherry picked from commit 8873a4c)
## Summary The TIR CSE pass currently lifts bool-typed sub-expressions like `i < n` or `a && b` into `cse_v: bool = ...` bindings whenever they appear twice. Boolean expressions are almost always predicates feeding `if` / `Select` / `assert`, where reading the condition inline is clearer than going through a boolean temporary, and where downstream simplification (ProveCondition, branch elimination) benefits from seeing the predicate directly. - Extend `CSEPlanner::IsEligible` in `src/tirx/transform/common_subexpr_elim.cc` to reject any compound expression whose result dtype is `bool`. - Update the file-level `Eligibility rules` doc-comment and the per-function `IsEligible` docstring to document the new rule. - Add two regression tests (`test_no_lift_bool_predicate`, `test_no_lift_bool_logical`) covering comparison predicates and logical-And predicates respectively. (cherry picked from commit e4c5b7c)
…d BATCH_TO_SPACE_ND (apache#19499) **Changes** Add tests in `test_frontend_tflite.py`. Lower S`PACE_TO_BATCH_ND` / `BATCH_TO_SPACE_ND` through TOPI in `tflite_frontend.py`. Use tf.raw_ops.BatchToSpaceND in the test because tf.batch_to_space_nd is not available in this TF build. **Why the TFLite frontend changed** The frontend was calling relax.op.nn.space_to_batch_nd / relax.op.nn.batch_to_space_nd, which aren’t implemented in this checkout. I updated the TFLite frontend to lower these ops via TOPI packed calls so conversion works and the new tests can pass. **Test:** ``` pytest test_frontend_tflite.py -k "test_space_to_batch_nd or test_batch_to_space_nd" ``` related to apache#18971 (cherry picked from commit 87bf302)
…pache#19497) `topi.scatter_elements` and `topi.scatter_nd` emit bare `T.parallel` loops in their te.extern IRBuilder bodies which trips `VerifyMemory` on CUDA targets: RuntimeError: Memory verification failed ... Did you forget to bind? CPU (LLVM) is unaffected. This fix makes the IRBuilder body in both `topi/scatter_elements.py` and `topi/scatter.py` target-aware. When `Target.current()` is a GPU target it emits thread bindings instead of `T.parallel`. Fixes apache#19451. (cherry picked from commit fde09d2)
…ape/Tensor inputs (apache#19498) ## Description When `from_onnx(model, keep_params_in_input=True)` is used, every ONNX initializer becomes a `relax.Var` instead of a `relax.Constant`. The `Concat` handler's `is_shape_like()` check only recognizes `relax.ShapeExpr` and 1D-int64 `relax.Constant`, so a 1D-int64 shape value loaded as a Var is no longer recognized. When such a Var is concatenated with a `ShapeExpr` — the standard pattern for dynamic-batch `Reshape` in PyTorch-exported ONNX models — the heterogeneous `Tuple(ShapeExpr, Tensor)` is rejected by `relax.op.concat` with: ``` InternalError: Op(relax.concat) expects the input to be a Tuple of Tensors. However, the given input is R.Tuple(R.Shape([N]), R.Tensor((1,), dtype="int64")) ``` This effectively breaks `keep_params_in_input=True` for any model with dynamic-batch `Reshape` (extremely common in PyTorch ONNX exports). ## Fix Run each `Concat` input through the existing `get_constant` helper before the `is_shape_like` check. This resolves any `Var` that maps to a known param back to its baked `Constant`, restoring the all-shape-like fast path. ## Minimal repro An 8-node ONNX graph (`Shape` → `Slice` → `Concat([dyn_n, [12]])` → `Reshape`) fails with `keep_params_in_input=True` before this PR and passes after. A regression test (`test_concat_with_param_shape_value`) covers this pattern. ## Testing ``` pytest tests/python/relax/test_frontend_onnx.py -k concat ``` 9 passed (1 new + 8 existing). (cherry picked from commit 82a37da)
Add OPFS as an alternative caching mechanism for artifacts: https://developer.mozilla.org/en-US/docs/Web/API/File_System_API/Origin_private_file_system. (cherry picked from commit 75a48a0)
…pache#19511) ## Motivation PyTorch's `torch.flip(x, dims=[...])` reverses every listed axis. The Relax converter `_flip` (`base_fx_graph_translator.py`) instead coerces the list to a single integer: ```python if isinstance(dims, list | tuple) and len(dims) > 0: dims = dims[0] ``` Only the first axis is forwarded to `relax.op.flip`, which is itself single-axis. The remaining axes are silently dropped. Minimal repro (vs PyTorch eager) on a `(3, 4)` input with `dims=[-1, -2]`: ``` ref: [11, 10, 9, 8, 7, 6, 5, 4, ...] # both axes flipped tvm: [ 3, 2, 1, 0, 7, 6, 5, 4, ...] # only last axis flipped ``` max_abs_diff = 8.0. Both the `torch.export` and legacy fx paths share this converter, so both are affected. ## Fix Iterate over `dims` in the converter and emit one `relax.op.flip` per axis (flips along distinct axes commute, so the order is irrelevant). A scalar `dims` is wrapped to a single-element list; non-int / non-sequence arguments still raise `TypeError`. `relax.op.flip` itself is unchanged: it is used elsewhere as a single-axis op, and widening its signature would expand the scope of this fix beyond the PyTorch frontend. (cherry picked from commit 61b49bb)
…e#19512) ## Motivation The PyTorch frontend's `_var` ignored the `correction` kwarg of `aten.var.correction`. `torch.export.run_decompositions()` rewrites both `aten.std.correction` and `aten.std.dim` into `aten.var.correction(..., correction=<value>) → sqrt`, so every `torch.std`/`torch.var` call lands in `_var` — but the correction value was dropped on the floor. The variance was therefore always divided by `n` regardless of what the user requested. Minimal repro (vs PyTorch eager): ``` x = [[1, 2, 3, 4, 5], [2, 2, 2, 2, 2]] torch.std(x, dim=1, unbiased=True) ref: [1.5811, 0.0] # sqrt(2.5) tvm: [1.4142, 0.0] # sqrt(2.0) -- correction silently set to 0 ``` The same omission shows up for explicit `torch.var(x, correction=k)` and any model that relies on the documented Bessel default. ## Fix Route `aten.var.correction` (identified by `OpOverload._overloadname`, not a substring match) to a new `_var_correction` helper. It reads `correction` from `node.kwargs`, treats `None` as 1 to match the overload's `Scalar? correction = None` schema, and scales the existing `relax.op.variance` output by `n / (n - correction)` when `correction != 0`. When `n - correction <= 0`, the multiplier is set to NaN rather than raising — this mirrors PyTorch's documented `max(0, N - correction)` semantics (eager produces NaN with a warning, not an error). Reduction-axis sizes are read from `x.struct_info.shape`. Dynamic sizes raise `NotImplementedError`; static-shape models cover the real-world `torch.export` flow. The legacy fx path through `_var` is intentionally left alone — it has a separate preexisting bug (it reads `args[2]` as `keepdim` even when that slot is `unbiased`), but fixing that here would expand the scope of this PR beyond the `correction` semantics. ## Notes - `_std` is also registered for `"std.correction"` but is unreachable on the default exported-program path because `aten.std.*` always decomposes to `var.correction + sqrt` before dispatch. Sparse-tensor exports that skip `run_decompositions` still hit the old `_std`; that path is out of scope for this fix. - Existing `test_std`/`test_var` encoded the buggy `correction=0` IR for `torch.var(x)` (which defaults to Bessel) and have been updated to expect the correct `R.multiply(var, R.const(15/14))`. New `test_var_correction` covers explicit `correction=2` and `correction=0`. (cherry picked from commit 5a7da7a)
…d root-block crash (apache#19514) ## Problem Closes apache#17873. `DefaultGPUSchedule` crashes when a PrimFunc body is a bare `SBlockRealize` (a fully-scalar op with no enclosing loops and no iter vars): ``` ValueError: Check failed: (sref->parent != nullptr) is false: Cannot add loops on top of the root block ``` Minimal repro (TVMScript decorators are omitted in this snippet to satisfy the PR-body lint; the regression test uses the regular `T.prim_func` form): ``` ir_module: prim_func main(a: Buffer((), "float32"), b: Buffer((), "float32"), c: Buffer((), "float32")): func_attr({"target": target("nvidia/geforce-rtx-3080")}) with sblock("scalar_add"): c[()] = a[()] + b[()] s_tir.transform.DefaultGPUSchedule()(M) # crashes ``` ## Root Cause The realized `scalar_add` block is itself the prim_func body's root sref — it has no parent stmt to mutate. `ThreadBind` (`src/s_tir/transform/default_gpu_schedule.cc`) reaches the `loops.empty()` branch and calls `sch->AddUnitLoop(block)`, which fails the `sref->parent != nullptr` check in `s_tir::AddUnitLoop` (`src/s_tir/schedule/primitive/loop_transformation.cc:1166`). The schedule infrastructure additionally requires the prim_func body to be an `SBlockRealize` whose block is the function's root (`GetRootPrimFunc` in `src/s_tir/schedule/analysis/analysis.cc:53`), so the body cannot simply be wrapped in a top-level `For`. ## Fix Before constructing the schedule, rewrite GPU-bound PrimFuncs whose body is a bare-leaf `SBlockRealize` so the realized block is no longer the root. The wrap conditions are intentionally narrow: 1. `func->body` is `SBlockRealize`, 2. the realized block has empty `iter_vars`, and 3. the block's body is not `For` or `SBlockRealize` (i.e. it is a leaf computation, not the well-formed implicit root that wraps a loop nest produced by the rest of the pipeline). When all three hold, the body becomes: ``` SBlockRealize( block=SBlock("root", body= For(u, 0, 1, kSerial, SBlockRealize(iter_values=[u], block=<original block, iter_vars=[IterVar(0..1, vu, kDataPar)]>)))) ``` The synthesised 1-extent data-parallel iter keeps `iter_values.size() == iter_vars.size()` for downstream checks, and the new For loop gives `ThreadBind` a real loop to bind to `blockIdx.x` / `threadIdx.x`. Already-scheduled functions and host-only PrimFuncs are skipped via the existing `IsScheduledOnGPU` / `kIsScheduled` gating. ## Testing ``` pytest tests/python/s_tir/transform/test_s_tir_transform_default_gpu_schedule.py ``` 10 passed (9 existing + 1 new `test_scalar_block_no_loops`). End-to-end compile + execute on RTX 3080 (sm_86): the scalar repro returns the expected `2.0 + 3.0 = 5.0`.
…9516) This adds explicit Expected IRModule coverage for TFLite GATHER and GATHER_ND frontend conversion. GATHER_ND uses Relax gather_nd with int64 indices, so the frontend now casts int32 TFLite indices to int64 before emitting the Relax op. This keeps the generated module well-typed and matches the expected Relax IR. Testing: - `python -m pytest tests/python/relax/test_frontend_tflite.py -k "gather"` related to apache#18971
…s index_put_ with tuple output (apache#19488) Hi Committers, This PR is trying to fix issues apache#18363. Any suggestions would be appreciated if you are available. ### Root Cause - When an ExportedProgram's FX graph output node returns a **nested Python tuple** (e.g., buffer mutation outputs + user-defined tuple returns), `_translate_fx_graph()` passes the raw nested structure directly to the Relax FFI Tuple constructor. - The C++ Array<Expr> initializer cannot handle heterogeneous/nested Python containers, causing a segmentation fault at `expr.cc`. - Additionally, index_put_ (in-place write op) did not update self.env to alias the source tensor to the mutated output, causing subsequent FX nodes that read the same tensor to observe **stale pre-mutation values**. ### Solution - exported_program_translator.py - Added static method `_flatten_output_args()` that recursively walks any Python `tuple/list`, collects only `relax.Expr` leaves, and preserve explicit None outputs as Relax null objects. - Replaced the fragile `assert isinstance(output_args, tuple | relax.Tuple)` guard with a call to `_flatten_output_args()`, producing a clean flat tuple of `relax.Expr` before FFI construction. - base_fx_graph_translator.py - In `_index_put()`, after emitting the `relax.op.index_put(...)` call, added an env alias update: `self.env[source_node] = output` when the target op name starts with `index_put_`, preserving correct in-place mutation semantics for downstream FX nodes. --------- Co-authored-by: cchung100m <cchung100m@users.noreply.github.com>
Description
This PR adds support for the CONV_3D operator in the TFLite frontend for
Relax.
Key Changes
- Operator Mapping: Added CONV_3D to the OperatorConverter mapping in
tflite_frontend.py.
- Implementation:
- Implemented convert_conv3d to handle 3D convolution attributes such as
StrideD/H/W, DilationD/H/W, and Padding.
- Correctly handled the TFLite 3D kernel layout, which is expected to be
DHWIO (Depth, Height, Width, Input Channels, Output Channels).
- Integrated support for fused activation functions (ReLU, ReLU6, etc.)
directly following the convolution.
- Unit Tests:
- Added comprehensive tests in
tests/python/relax/test_frontend_tflite.py covering:
- VALID and SAME padding modes.
- Various stride and dilation configurations.
- Verification against expected Relax IR structure.
Testing:
- `python3 -m pytest tests/python/relax/test_frontend_tflite.py -k
"test_conv3d"`
Notes for Reviewers
The implementation follows the existing pattern used for CONV_2D but
extends it to the 5D case (NDHWC layout). I've ensured that the kernel
layout mapping aligns with TVM's R.nn.conv3d requirements.
Related to: apache#19519
## Summary `AttrFunctor` is declared infrastructure with no remaining users. An exhaustive search across `include/`, `src/`, `tests/`, `python/`, `apps/`, `web/`, and `cmake/` confirms zero subclasses, zero friend declarations, and zero macro callers outside the header itself — its two internal macros (`ATTR_FUNCTOR_DEFAULT`, `ATTR_FUNCTOR_DISPATCH`) are only used inside `src/ir/attr_functor.h`. The one `#include "attr_functor.h"` in `src/ir/attrs.cc` is a stale leftover from a prior migration; that file references only `DictAttrs` and `AttrFieldInfoNode` (from `tvm/ir/attrs.h`), not any `AttrFunctor` symbols. Removes `src/ir/attr_functor.h` (150 lines) and drops the stale include. No build-system changes needed — the header was never enumerated in any `CMakeLists.txt`.
…ther` operator (apache#19525) Hi Committers, This PR is trying to fix issues apache#19436. Any suggestions would be appreciated if you are available. ### Root Cause 1. ONNX `Gather` allows negative indices (counting from the end of the target axis). 2. In the Relax ONNX importer, `Gather` was lowered directly to `relax.op.take` without normalizing negative indices first. 3. This created semantic mismatch / incorrect behavior in downstream lowering paths that assume non-negative indices. 4. Test failures were also caused by pytest parametrization issues: - using ONNX `TensorProto` enum values directly as NumPy dtypes, - and tuple-style parametrization triggering fixture interpretation errors. ### Solutions 1. Added conditional negative-index normalization in `Gather._impl_v13`: - apply only for signed index dtypes, - use: `idx < 0 ? idx + axis_extent : idx`, - derive `axis_extent` from shape/runtime expression to support dynamic shapes. 2. Skipped normalization for unsigned index dtypes to avoid redundant graph ops/checks. --------- Co-authored-by: cchung100m <cchung100m@users.noreply.github.com>
…pache#19530) This commit adds support for the CONV_3D_TRANSPOSE operator in the Relax TFLite frontend. Key implementations: - Registered CONV_3D_TRANSPOSE to the TFLite op map. - Implemented convert_conv3d_transpose which shares Conv3DOptions with regular Conv3D but handles the distinct tensor input layout [output_shape, weight, data, bias] and the DHWOI kernel layout. - Added calculation for SAME padding that correctly handles transposed convolution semantics, computing padding and output_padding based on dilated kernel and stride sizes. - Added comprehensive unit tests for valid and same padding in test_frontend_tflite.py. Testing: - `python3 -m pytest tests/python/relax/test_frontend_tflite.py -k "test_conv3d_transpose"` Related to: apache#19519
…orage scope (apache#19423) part of tile-ai/tilelang#1869 ## Summary Add TIR builtins and storage scope for Metal cooperative_tensor operations (MetalPerformancePrimitives / Metal 4). ## Motivation Apple Metal 4 introduces MetalPerformancePrimitives (MPP) with `matmul2d` using `cooperative_tensor` operands. On M5, this routes to NAX tensor cores; on M1-M4, it falls back to simdgroup matrix instructions. These TIR primitives enable backend codegen to emit MPP calls. ## Changes ### New TIR builtins - `cooperative_tensor_fill(d, index, value, rows, cols)` - `cooperative_tensor_load(d, index, ptr, stride, rows, cols, transpose)` - `cooperative_tensor_store(d, index, ptr, stride, rows, cols, transpose)` - `cooperative_tensor_multiply_accumulate(d, di, a, ai, b, bi, c, ci, M, N, K, trans_a, trans_b)` ### New storage scope - `metal.cooperative_tensor` (`StorageRank::kMetalCooperativeTensor`) ### Files changed - `include/tvm/tirx/builtin.h` — Op declarations - `src/tirx/op/builtin.cc` — Op registrations - `python/tvm/tirx/op.py` — Python wrappers - `python/tvm/script/ir_builder/tirx/ir.py` — Script parser exports - `src/runtime/thread_storage_scope.h` — StorageRank enum + scope parsing These builtins mirror the existing `simdgroup_*` builtins for the older Metal simdgroup matrix API, extended with M/N/K dimension parameters for the matmul2d descriptor.
…ort (apache#19536) ## Summary This PR adds initial Relax TFLite frontend support for 29 StableHLO builtin operators from apache#19519 item I. The covered subset includes pure elementwise ops, BuiltinOptions2 / metadata-based ops, simple shape-manipulation ops, and a take-equivalent subset of `STABLEHLO_GATHER`. StableHLO builtins carry no TFLite-specific quantization or fused-activation metadata, so the implementation uses dedicated converter helpers that bypass the existing TFLite elemwise/QNN code paths. Relates to apache#19519. ## Changes 1. **Zero-attribute elementwise helpers** - Add `_convert_stablehlo_unary`, `_convert_stablehlo_binary`, and `_convert_stablehlo_ternary` for pure elementwise mapping. - Register 20 ops: unary (`ABS`, `NEGATE`, `COSINE`, `EXPONENTIAL`, `FLOOR`, `LOG`, `LOGISTIC`, `RSQRT`, `TANH`), binary (`ADD`, `SUBTRACT`, `MULTIPLY`, `DIVIDE`, `MAXIMUM`, `MINIMUM`, `POWER`), ternary (`SELECT` → `R.where`), and dtype-dispatched bitwise/logical ops (`AND` / `OR` → logical ops for bool or bitwise ops for integer, `SHIFT_LEFT` → `R.left_shift` for integer). 2. **BuiltinOptions2 infrastructure** - Add `_get_stablehlo_options` helper for parsing `BuiltinOptions2` flatbuffers with enum validation via `getattr(BuiltinOptions2, options_cls.__name__)`. - Register 6 ops: `CONVERT` → `R.astype`, `CLAMP` → `R.minimum(R.maximum(...))`, `CONCATENATE` → `R.concat`, `BROADCAST_IN_DIM` → `R.reshape` + `R.broadcast_to`, `IOTA` → `R.arange` + `R.broadcast_to`, and `COMPARE` → 6 comparison directions (`TOTALORDER` raises `OpNotImplemented`). 3. **Shape-manipulation ops** - `PAD` → `R.nn.pad` in constant mode. The initial PAD path supports non-negative edge padding with zero interior padding and a constant scalar padding value. Interior padding, negative padding, and dynamic padding values raise `OpNotImplemented`. - `DYNAMIC_SLICE` → `R.dynamic_strided_slice`. The initial path supports constant, in-bound start indices only. Runtime start indices and out-of-bounds StableHLO clamping semantics are deferred. 4. **Indexing op** - `GATHER` → `R.take` for the take-equivalent subset only. - Parses the relevant `StablehloGatherOptions` attributes needed to validate this subset: `offset_dims`, `collapsed_slice_dims`, `start_index_map`, `index_vector_dim`, and `slice_sizes`. - Validates the gather axis, collapsed dims, offset dims, slice sizes, and output shape against the expected `R.take` layout. Multi-dimensional and non-take-equivalent gather patterns raise `OpNotImplemented`. 5. **Not included** - `STABLEHLO_RESHAPE`, `STABLEHLO_TRANSPOSE`, and `STABLEHLO_SLICE` are left to another contributor who expressed interest in those ops. - The remaining Issue apache#19519 StableHLO items are deferred to follow-up PRs: `CBRT`, `REMAINDER`, `SCATTER`, `CONVOLUTION`, `DOT_GENERAL`, `REDUCE`, `REDUCE_WINDOW`, `DYNAMIC_UPDATE_SLICE`, `COMPOSITE`, `CUSTOM_CALL`, `RNG_BIT_GENERATOR`, `SORT`, and `WHILE`. - More general or multi-dimensional `STABLEHLO_GATHER` patterns are also deferred to follow-up work. ## Testing All tests use manually-built minimal TFLite flatbuffers with `tvm.ir.assert_structural_equal`. BuiltinOptions2 ops construct their options via the FlatBuffers schema API, modeled after the existing DILATE test pattern. ```bash python -m pytest tests/python/relax/test_frontend_tflite.py -k stablehlo -q ``` ## Result - 29 StableHLO operators registered in the Relax TFLite frontend. - 44 StableHLO test cases covering all registered ops, including structural-equal tests and unsupported/error-path checks: - `COMPARE` with `TOTALORDER` - `PAD` with interior padding, negative padding, and dynamic padding values - `DYNAMIC_SLICE` with runtime starts and out-of-bounds starts - non-take-equivalent or multi-dimensional `GATHER` - All StableHLO TFLite frontend tests pass locally. ## References - Issue apache#19519 item I: StableHLO operators in TFLite - Related PR apache#19481: DILATE operator mapping, the first use of BuiltinOptions2 in the TFLite frontend tests
) Six CUDA sources in src/runtime/contrib used LOG(FATAL) via transitive includes that apache#19483 trimmed; add the explicit <tvm/runtime/logging.h> include to thrust.cu, attention_kernels.cu, and the four cutlass kernel headers (fp16/fp8 sm90/sm100, gemm_runner, fp8_groupwise_scaled_gemm). cache_kernels.cu used the bare Array{...} alias that apache#19483 removed; switch to ffi::Array<Tensor>{...}. attention_kernels.cu registered FFI functions whose parameters were raw DLTensor*; the new reflection registry requires TypeSchema, so wrap both TVM_FFI_STATIC_INIT_BLOCK registrations to take Tensor and forward to the unchanged launchers via GetDLTensorPtr() (with const_cast for the output tensors, matching the mt_random_engine / cudnn pattern).
### Summary - Respect the ONNX `reduction` attribute in the Relax ONNX frontend `ScatterElements` converter. - Preserve existing default behavior by mapping missing reduction and ONNX `none` to Relax `update`. - Add focused regression coverage for opset 11 default behavior, opset 16 `add`/`mul`, and opset 18 `none`/`min`/`max`. ### Changes - Added a shared helper to normalize and validate ONNX reduction attributes. - Implemented `ScatterElements` opset 16 and opset 18 converters. - Reused the existing `relax.op.scatter_elements(..., reduction=...)` API. - Reused the same reduction helper in `ScatterND` to keep behavior consistent. ### Test Plan - `python -m py_compile python/tvm/relax/frontend/onnx/onnx_frontend.py tests/python/relax/test_frontend_onnx.py` - `python -m pytest tests/python/relax/test_frontend_onnx.py::test_gather_elements tests/python/relax/test_frontend_onnx.py::test_scatter tests/python/relax/test_frontend_onnx.py::test_scatter_elements_reduction tests/python/relax/test_frontend_onnx.py::test_scatter_nd -q` ### Issue Fixes apache#19435 ## Local Verification Notes - WSL conda environment: `/home/thinker/.cache/tvm-conda-onnx` - TVM build directory: `/home/thinker/.cache/tvm-build-onnx` - LLVM runtime check: `tvm.runtime.enabled("llvm") == True` - Relevant ONNX frontend subset: `15 passed, 4 skipped, 2 warnings` - Full `tests/python/relax/test_frontend_onnx.py` was also attempted. It currently has 14 failures in unrelated `Reduce* axes input` and `TopK` tests; running the same selected failures against `origin/main` reproduces them, so they are not introduced by this PR.
apache#19535) This PR fixes apache#19533: - Sanitize floating tensor min/max: replace NaN with +inf/-inf before topi max/min so bounds match ONNX "unbounded" semantics where NaN bounds default to no constraint. - After clamping, preserve NaNs from the input tensor on floating dtypes. - Extend check_correctness with equal_nan for float outputs containing NaN. - Add parametrized Clip opset-13 tests for NaN min/max tensor bounds.
This PR fixes apache#19552. astral-sh/setup-uv is not on the ASF GitHub Enterprise action allowlist, causing the Lint workflow to fail with "Startup failure" before any pre-commit checks run. See https://github.com/apache/tvm/actions/runs/25743684906 for the failed reason. This PR removes the uv setup and sync steps entirely; pre-commit/action will install and manage pre-commit and all hook dependencies on its own. This PR also corrected previous lint errors. After the fix, the CI lint succeeded: https://github.com/apache/tvm/actions/runs/25775499703/job/75707088129
… NonMaxSuppression (apache#19547) Hi Committers, This PR is trying to fix issues apache#19544. Any suggestions would be appreciated if you are available. --------- Co-authored-by: cchung100m <cchung100m@users.noreply.github.com>
…apache#19515) ## Summary Introduce a test runner that reuses the official ONNX Backend Test Suite to systematically verify the Relax ONNX importer. This complements the existing hand-written tests in `test_frontend_onnx.py` by providing spec-aligned coverage of standard ONNX operator semantics. Towards apache#19505 ## Motivation The existing `test_frontend_onnx.py` has 187 hand-written tests that validate TVM-specific importer behavior (parameter handling, name sanitization, dynamic shapes, Relax IR structure). However, it relies on ONNX Runtime as the reference and cannot systematically cover all edge cases defined in the ONNX specification. The ONNX Backend Test Suite provides 1653+ node-level tests with protobuf reference inputs/outputs. It is the industry standard for validating ONNX importers/exporters (used by ONNX Runtime, TensorFlow, PyTorch). Reusing it gives Relax a living, upstream-aligned correctness baseline. ## What this PR adds - `tests/python/relax/test_frontend_onnx_backend.py` — a backend adapter (`TVMRelaxBackend`) that implements the `onnx.backend.base.Backend` interface, wiring `from_onnx()` → `DecomposeOpsForInference()` → `LegalizeOps()` → `tvm.compile()` → `VirtualMachine`. ## Coverage 72 operators with 388 test cases, all passing. Only operators where every ONNX node test passes are included — no xfail markers. Operators not yet covered include: cast (exotic dtypes), reduce ops (edge cases), reshape/resize/attention (complex behavior), quantization, and several others with known importer gaps. These can be added incrementally as the importer improves. ## Test results 388 passed, 3216 skipped (CUDA variants + operators not yet in allowlist), 0 failed, 0 xfailed ## CI impact - New test file is not added to any existing CI test shard by default - Full suite (388 tests) is lightweight on CPU-only runners ## Design decisions - **Coexistence with existing tests**: `test_frontend_onnx.py` remains unchanged. Backend tests cover standard ONNX semantics; hand-written tests continue to cover TVM-specific behavior (dynamic shapes, Relax IR structure, importer options). - **Public API only**: uses `backend_test.include()` with `^`-anchored regex patterns. No access to private ONNX APIs. - **No xfail**: only include operators that fully pass. Uncovered operators are documented in code comments and this PR description. Follow-up PRs can expand coverage as importer gaps are fixed. - **Prefix conflict handling**: `include()` patterns use `^test_{op}(?:_.*)?(?:_cpu|_cuda)$`, which can cause false matches when a short op name is a prefix of a longer one (e.g. `log` vs `log_softmax`). Affected ops (`log`, `max`, `relu`) are excluded until a more precise matching strategy is adopted.
This PR fixed apache#19551. Bool product has logical-AND semantics and cannot be lowered through TIR Mul for LLVM codegen. Route bool prod through all() and add frontend and legalization coverage for bool R.prod.
Hi Committers, This PR is trying to fix issues apache#19541. Any suggestions would be appreciated if you are available. ### Root cause: The ONNX `Div` path in the Relax frontend did not separate two integer-divisor cases: constant zero divisors and dynamic/unknown divisors. As a result, constant integer zero divisors were not rejected during import, and dynamic integer divisors could reach runtime without a guard. When the divisor became zero at runtime, execution could trigger SIGFPE and terminate the process instead of raising a controlled error. ### Solution: This PR applies a minimal, targeted fix in the ONNX frontend `Div` conversion path. It introduces: import-time validation for constant integer divisors containing zero, raising ValueError early. --------- Co-authored-by: cchung100m <cchung100m@users.noreply.github.com>
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. - 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. - `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` 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.
… Taylor (apache#19567) ## Summary `tirx.asin`'s LLVM legalize used a 6-term Taylor series for `|x| < 0.5` with wrong recurrence coefficients. The ratios in the code (`9/40`, `25/112`, `1225/3456`, `3969/28160`) don't match the real asin series (`9/20`, `25/42`, `49/72`, `81/110`), so mid-range inputs lose ~1e-3 of precision — over 1000 float32 ULP. `acos` inherits it via `π/2 − asin(x)`. ``` x=0.47 ORT=0.48929077 TVM(old)=0.48820966 err=-1.08e-3 ``` The Taylor branch was added in apache#17945 as the initial implementation, with no libm fallback. apache#18582 later patched only `|x| ≥ 0.5` by routing to the libm extern, leaving the buggy mid-range in place. I see no evidence the inline series was an intentional fast-path. ## Fix Drop the inline series, route the whole domain through the existing `asinf`/`acosf` extern, keep the out-of-range NaN guard. Max error over `x ∈ [-1, 1]` drops to **2.4e-7** (ULP-grade). ## Tests - Re-enable `Asin`/`Acos` in `test_unary` (they were commented out with a TODO about Taylor precision loss). - Existing `test_asin_acos_boundary_values` (apache#18582) still passes. If the inline polynomial was intentional for some target/path, please flag it — I'll restore it with corrected coefficients instead. `Atan` is still disabled; that's a separate `x² + 1` overflow bug (apache#19560). Fixes apache#19563.
… option (apache#19565) Fix CUDA lowering of standard TIR math intrinsics so they use precise CUDA math functions by default instead of fast-math `__*f` functions. This fixes the default behavior reported in apache#19546, where operators such as `tirx.exp` could lower to `__expf` even though fast math was not explicitly requested. This change adds a CUDA target attribute, `enable_fast_math`, which defaults to `false`. When the attribute is unset or false, standard math intrinsics lower through the normal CUDA math rule, for example `expf`, `logf`, `sinf`, `cosf`, `powf`, and `rsqrtf` for `float32`. When users explicitly enable the attribute on the target, the lowering pass also checks the `cuda.fastmath.FLowerIntrinsic` rules before the normal CUDA lowering rules. Users can opt in to fast math by constructing a CUDA target with the attribute: ```py tvm.target.Target({"kind": "cuda", "enable_fast_math": True}) target = tvm.target.Target({ "tag": "nvidia/nvidia-a100", "enable_fast_math": True, }) ``` The fast-math lowering path currently covers the CUDA math operators registered with `cuda.fastmath.FLowerIntrinsic`: `tirx.exp`, `tirx.exp10`, `tirx.log`, `tirx.log2`, `tirx.log10`, `tirx.tan`, `tirx.cos`, `tirx.sin`, `tirx.tanh`, and `tirx.pow`. `tirx.rsqrt` is also registered for CUDA lowering so it maps to the CUDA reciprocal-square-root intrinsic instead of being legalized as `1 / sqrt(x)`. Add CUDA codegen tests `tests/python/codegen/test_target_codegen_cuda_fastmath.py` that check the lowered IR, generated CUDA source, and runtime results for the supported math intrinsics across floating point dtypes and both default and fast-math targets.
…on (apache#19643) This PR will fix apache#19592. LayerNorm could produce NaN on large-value, small-variance inputs due to catastrophic cancellation in var = E[x^2] - E[x]^2. Switch to a numerically stable two-pass formulation: - pass1 computes mean via sum(x) / N - pass2 computes variance via sum((x - mean)^2) / N (cherry picked from commit fca86a6)
… same tensor (apache#19644) This PR will fix apache#19577. In this issue, the IRModule before applying any pass looks like: ``` %x: Tensor[(4,), float32] // function param with R.dataflow(): %lv = expand_dims(%x, axis=1) // (4, 1) %lv1 = expand_dims(%x, axis=1) // (4, 1) second call, new Var %lv2 = multiply(%lv, %lv1) // (4, 1) %lv3 = concat(%lv2, %lv1, axis=1) // (4, 2) ... ``` When the users manually apply the `DataflowUseInplaceCalls` pass, the pass will rewrite the statement `%lv2 = multiply(%lv, %lv1)` to be like `%lv = multiply(%lv, %lv1); %lv3 = concat(%lv, %lv1, axis=1)`, which reuses the %lv buffer to avoid storage waste. But this rewrite will chang the buffer context of %lv, and also in LLVM generated code, %lv1 shared the same storage with %lv, so when executing `%lv = concat(%lv, %lv1, axis=1)`, the %lv1 context has also been changed to `multiply(%lv, %lv1)`. So the failure is due to the shared storage of different views of the same tensor %x. During the execution, %lv1 holds `x^2` instead of `x` after `multiply`. `concat` reads %lv1 for the right column and its result is [[1,1],[4,4],[9,9],[16,16]] instead of [[1,1],[4,2],[9,3],[16,4]] (the correct result should be : left col `x^2`, right col should stay `x`). Change: View-like ops (expand_dims, squeeze, reshape, permute_dims, memory.view, ensure_zero_offset) take the input's alias set in alias analysis instead of a new id: %lv and %lv1 share alias with %x. Then the pass rejects in-place of `multiply(%lv, %lv1)`: %lv and %lv1 are different vars but alias ids intersect, so no operand may be reused in-place. (cherry picked from commit d705cb2)
## Summary This PR adds conservative Relax TFLite frontend support for the TFLite builtin `STABLEHLO_CUSTOM_CALL` operator. TFLite marks `STABLEHLO_CUSTOM_CALL` as having no runtime kernel. Importing general custom calls as executable Relax operators would therefore give them semantics that TFLite itself does not provide. This PR only supports the metadata-only `Sharding` custom call target, which TensorFlow's StableHLO pipeline treats as an annotation that can be erased. ## Design ### Sharding Annotation Lowering `STABLEHLO_CUSTOM_CALL` now parses `StablehloCustomCallOptions` from `BuiltinOptions2` and reads the `call_target_name`. For `call_target_name == "Sharding"`, the frontend lowers the op to identity: the output tensor is bound to the input expression. This mirrors TensorFlow's handling of Sharding custom calls as metadata annotations. The sharding spec in `backend_config` is intentionally dropped for single-device import. The supported subset is guarded: - exactly one input and one output - input and output shape/dtype metadata must match - `has_side_effect` must be false - `called_computations` must be empty All other custom-call targets raise `OpNotImplemented` with the target name in the diagnostic. ## Operator Support | Operator | TFLite options | Relax lowering | Supported subset | |---|---|---|---| | `STABLEHLO_CUSTOM_CALL` | `StablehloCustomCallOptions` from `BuiltinOptions2` | identity for `Sharding`; otherwise unsupported | metadata-only `Sharding` annotations with unchanged tensor metadata | ## Tests The tests manually build minimal StableHLO custom-call TFLite flatbuffers and compare the supported identity path with `tvm.ir.assert_structural_equal`. Unsupported patterns use `pytest.raises`. | Test | Coverage | |---|---| | `test_stablehlo_custom_call_sharding` | `Sharding` annotation lowers to identity | | `test_stablehlo_custom_call_unsupported_target` | unknown external target guard | | `test_stablehlo_custom_call_sharding_side_effect_unsupported` | side-effecting `Sharding` guard | | `test_stablehlo_custom_call_sharding_metadata_mismatch_unsupported` | input/output metadata guard | Local validation: ```bash python -m py_compile \ python/tvm/relax/frontend/tflite/tflite_frontend.py \ tests/python/relax/test_frontend_tflite.py python -m ruff check \ python/tvm/relax/frontend/tflite/tflite_frontend.py \ tests/python/relax/test_frontend_tflite.py python -m pytest \ tests/python/relax/test_frontend_tflite.py \ -k stablehlo_custom_call -q python -m pytest \ tests/python/relax/test_frontend_tflite.py \ -k stablehlo -q ``` Result: ```text py_compile: passed ruff check: All checks passed stablehlo_custom_call tests: 4 passed stablehlo tests: 81 passed ``` ## References - Issue apache#19519 item I: remaining StableHLO operators in TFLite - TensorFlow Lite schema marks `STABLEHLO_CUSTOM_CALL` as no runtime support - TensorFlow StableHLO pipeline erases `Sharding` custom calls as metadata annotations (cherry picked from commit cf859b9)
…pache#19653) In continuation of apache#19624 this catches some unlifted entries. Hope there is no more left, for consistency it now covers comments and perhaps non-active (hotpath) parts. (cherry picked from commit 349225a)
## Summary Add Relax TFLite frontend support for `HASHTABLE_LOOKUP`. This PR adds a converter for `HASHTABLE_LOOKUP` in the Relax TFLite frontend. The implementation supports non-string value tensors and lowers the lookup through `bucketize`, `take`, and `where` so that missing keys return zero-filled values together with a `uint8` hits mask matching TFLite semantics for the supported cases. The PR also adds handcrafted TFLite frontend tests covering: - 1D float value tensors - 2D float value tensors - the current unsupported string-value case ## Testing Ran `tests/python/relax/test_frontend_tflite.py -k 'hashtable_lookup'`. Part of apache#19519 (cherry picked from commit 066bf77)
…#19651) ## Summary This PR adds Relax TFLite frontend support for the TFLite builtin `STABLEHLO_RNG_BIT_GENERATOR` operator. Unlike most StableHLO builtins, the TFLite runtime (`tensorflow/lite/kernels/rng_bit_generator.cc`) implements this op as a real, deterministic counter-based PRNG, so the importer must reproduce it bit-exactly rather than map it to an existing op: - one uint64 1-D `initial_state` input, two outputs — uint64 `output_state` and the random-bit `output` (int32 / int64 / uint32 / uint64); - `algorithm` in `{DEFAULT, PHILOX, THREEFRY}`, where `DEFAULT` resolves to `PHILOX`; - Random123 Threefry2x32 (20 rounds) and Philox4x32 (10 rounds) with the fixed constants from `rng_util.cc`; - state-length constraints: `THREEFRY` requires `u64[2]`, `PHILOX`/`DEFAULT` require `u64[2]` or `u64[3]`. ## Design TVM/Relax has no matching RNG primitive, so the converter generates a TIR kernel that mirrors the runtime and emits it through `relax.call_tir` with two outputs. The kernel: - reinterprets the uint64 state as uint32 words and advances a 64-bit block counter (`final counter = initial_state[1] + num_blocks`); - runs the selected algorithm per block with all round state materialized into local buffers, which keeps the generated IR linear instead of an exponentially nested expression tree; - packs the produced uint32 words back into the output dtype, and writes the updated state (key unchanged, counter advanced, Philox `u64[3]` tail passed through) — the only state behaviour the runtime relies on. The kernel is an `s_tir` PrimFunc wrapped in a single opaque structured block so it remains a well-formed block-structured function for the Relax pipeline (e.g. `HasReshapePattern`). `get_tensor_type_str` and the input `_decode_type` map are extended with uint32/uint64 so the uint64 state imports correctly. Unsupported inputs raise a precise `OpNotImplemented` (non-uint64 / non-1-D state, mismatched output-state shape, unsupported output dtype, unknown algorithm, per-algorithm state-length violations). ## Operator Support | Operator | TFLite options | Relax lowering | Supported subset | |---|---|---|---| | `STABLEHLO_RNG_BIT_GENERATOR` | `StablehloRngBitGeneratorOptions.Algorithm()` from `BuiltinOptions2` | `call_tir` to a generated bit-exact TIR kernel | THREEFRY (`u64[2]`) and PHILOX/DEFAULT (`u64[2]`/`u64[3]`); int32/int64/uint32/uint64 output | ## Tests Tests build minimal RNG flatbuffers, compile, and execute them, comparing the output and updated state against the verbatim expected vectors from the TFLite runtime kernel test (`rng_bit_generator_test.cc`). | Test | Coverage | |---|---| | `test_stablehlo_rng_bit_generator_threefry` | THREEFRY bit-exact, all 4 output dtypes | | `test_stablehlo_rng_bit_generator_philox` | PHILOX bit-exact, all 4 output dtypes | | `test_stablehlo_rng_bit_generator_default_matches_philox` | DEFAULT resolves to PHILOX | | `test_stablehlo_rng_bit_generator_deterministic` | run-to-run bit-identical output | | `test_stablehlo_rng_bit_generator_unsupported_output_dtype` | output dtype guard | | `test_stablehlo_rng_bit_generator_threefry_invalid_state_unsupported` | THREEFRY `u64[2]` state guard | | `test_stablehlo_rng_bit_generator_non_uint64_state_unsupported` | uint64 state guard | Local validation: ```bash python -m ruff check \ python/tvm/relax/frontend/tflite/tflite_frontend.py \ tests/python/relax/test_frontend_tflite.py python -m pytest \ tests/python/relax/test_frontend_tflite.py \ -k rng_bit_generator -q python -m pytest \ tests/python/relax/test_frontend_tflite.py \ -k stablehlo -q ``` Result: ```text ruff check: All checks passed rng_bit_generator tests: 13 passed stablehlo tests: 96 passed ``` ## References - Issue apache#19519 item I: remaining StableHLO operators in TFLite - `tensorflow/lite/kernels/rng_bit_generator.cc`, `rng_util.cc`, `rng_bit_generator_test.cc` (cherry picked from commit 23a0ea8)
…ache#19648) ## Summary This patch resolves a security vulnerability by explicitly setting the `android:exported` flag within the `AndroidManifest.xml` file. Previously, certain components were missing this required flag, which could lead to incorrect permission handling and exposure, potentially allowing unauthorized access to components of the application. ## Changes * **Security:** Added explicit `android:exported="true"` or `android:exported="false"` flags to relevant `<activity>`, `<receiver>`, and `<provider>` tags in `AndroidManifest.xml`. * **Safety:** Ensures that all exposed components properly define their export status, adhering to modern Android best practices and mitigating potential misconfigurations. * **Compatibility:** Improves the application's security posture and adherence to Android framework requirements regarding component visibility. ## Testing - Verified logic locally using Docker sandbox Fixes #AUTO_SEMGREP Co-authored-by: CodeMechanic <codemechanic@local.ai> (cherry picked from commit 5649122)
apache#19645) ### Motivation `torch.logical_not` accepts an input tensor of any dtype (treating any nonzero element as `True`) and always returns a `bool` tensor. The PyTorch frontend previously lowered it with `self._unary_op(relax.op.logical_not)`. `relax.op.logical_not` is a unary arithmetic op that passes its input dtype through, so a non-bool input (for example `float32`) produced a `float32` result instead of the `bool` result PyTorch returns. This is a dtype mismatch against the reference PyTorch semantics for both the FX and ExportedProgram frontends. ### Changes - Add a shared `_logical_not` converter in `BaseFXGraphImporter` that casts non-bool inputs to `bool` before applying `relax.op.logical_not`. Bool inputs are passed through unchanged (no redundant cast). - Point the `logical_not` (FX) and `logical_not.default` (ExportedProgram) registrations at the new converter. - Update the FX test and add a standalone ExportedProgram `test_logical_not` to assert the corrected IR (`astype` to bool, then `logical_not`, 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 9898909)
The CrossOriginStorage class was storing the URL→hash map only in the module-level GLOBAL_HASH_CACHE. After a page reload that cache is empty, and getFileHash() can only recover hashes for HuggingFace LFS files (URLs containing /resolve/). This left several resource categories uncacheable across sessions: <img width="1197" height="279" alt="Screenshot 2026-05-15 at 17 43 15" src="https://github.com/user-attachments/assets/c9943910-9002-4b06-afdd-6288b7e22ba6" /> - JSON files not stored in LFS (mlc-chat-config.json, tokenizer.json, tensor-cache.json) — getFileHash returns null for their /resolve/ URLs because the raw pointer is the actual file content, not an LFS pointer. - .wasm files from GitHub raw URLs — no /resolve/ pattern at all. - Any file whose hash was computed from blob content via getBlobHash. Additionally, even for genuine LFS model shards, each page load was re-fetching every shard's LFS pointer file over the network just to re-derive the SHA-256 hash. Fix: persist the URL→hash mapping to a dedicated Cache API store (tvmjs-cos-hash-meta). Two write sites: 1. put() — after a file is stored in COS, persist its blob-derived hash. This covers all non-LFS files and non-HuggingFace URLs. 2. resolveHashDescriptor() — after getFileHash() resolves a hash from the LFS pointer, persist it immediately. This eliminates repeated pointer-file network requests for model shards on subsequent visits. Both write sites use a best-effort try/catch so storage quota errors are silently ignored. loadPersistedHashEntry() similarly swallows errors. The typeof caches === "undefined" guard keeps the code safe in Node.js test environments. (cherry picked from commit 8039963)
…ss (apache#19650) Fix a crash (apache#19576) when AdjustMatmulOrder encounters mixed-dimension matmul chains common in transformer models (e.g. matmul(attn_output[B,S,D], W_o[D,D])). The pass previously assumed all operands in a chained rewrite were 2D and asserted shape_c.size() == 2, failing on 3D intermediate results. Changes: - Replace full 2D transpose with permute_last_two_dims for permuted matmul patterns, swapping only the last two axes for ND tensors. - Remove hard ndim==2 checks in the permuted rewrite path. - Account for batch prefixes when comparing naive matmul FLOPs, so reorder decisions reflect batched vs. weight-only inner matmuls. - Skip reorder when neither evaluation order is provably cheaper. - Add regression tests for symbolic/concrete batched LoRA shapes. - Add a numerics test covering a minimal attention block with ND permute_dims. (cherry picked from commit dda158c)
…he#19652) ## Summary Add Relax TFLite frontend support for `EMBEDDING_LOOKUP_SPARSE`. This PR adds a converter for `EMBEDDING_LOOKUP_SPARSE` in the Relax TFLite frontend. The implementation supports the `SUM`, `MEAN`, and `SQRTN` combiners and handles higher-rank sparse indices. The sparse aggregation is lowered through `scatter_nd` to match TFLite operator semantics for the supported cases. The PR also adds handcrafted TFLite frontend tests covering: - `SUM` - `MEAN` - `SQRTN` - a 3D indices case ## Testing Ran `tests/python/relax/test_frontend_tflite.py -k 'embedding_lookup_sparse'`. Part of apache#19519 --------- Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com> (cherry picked from commit 4a688dd)
Add a workflow_dispatch pipeline that builds, tests, and publishes manylinux/ macOS/Windows wheels via cibuildwheel with OIDC trusted publishing. (cherry picked from commit 5cbf506)
…che#19657) ## Summary Follow-up work on top of the TIRx infrastructure bring-up (apache#19581). It extends the TIRx operator-dispatch, codegen, and TVMScript surfaces with the next batch of low-level programming features for Blackwell-class GPUs, while keeping `s_tir` script support intact. ## Main Changes - **op-dispatch**: warp `ldmatrix`/`stmatrix` copy dispatch; split CUDA copy into register / gmem-smem / `ldgsts` paths; `tcgen05.ld/st` `.16x{64,128,256}b` dispatch with a factory and M=128 layout; element-wise broadcast at the layout level with a copy vec-alignment fix. - **gemm**: CUDA synchronous `mma.sync` tensor-core dispatch; accept a Layout F C operand for M=64 MMAs. - **op**: add the `permute_layout` primitive (replaces `permute_dims`). - **tvmscript**: add the `Tx.jit` decorator, `Tx.constexpr` compile-time params, and `Tx.wg_reg_tile`. - **lower-tirx**: introduce the `Tx.device_entry()` marker (replacing `ScopeKind::kKernel`); canonical thread filters that drop the `Tx.filter` wrapper. - **codegen**: add a typed-pointer byte-offset intrinsic; remove the `entry_cluster_sync` codegen attribute. ## Validation - `pre-commit run` (changed files) — clean - `ninja -C build -j$(nproc)` — builds - `pytest tests/python/tirx/ -n 16` - `1997 passed, 39 skipped, 3 xpassed` - `python -m pytest tests/python/all-platform-minimal-test` - `37 passed, 105 skipped` - `TVM_TEST_TARGETS=llvm pytest tests/python/tirx-analysis tests/python/tirx-base tests/python/tirx-transform -n 16` - `630 passed, 25 skipped, 8 xfailed, 1 xpassed` ## Local CI Notes Several full CI-equivalent jobs are not locally reproducible because this machine is missing parts of the Apache TVM CI environment (e.g., specific `llvm-config` versions, Vulkan, ROCm, ARM/QEMU cross-toolchain, and web/wasm components). The Blackwell/Trainium kernel tests are maintained downstream and are intentionally not part of this PR. (cherry picked from commit 57c638f)
Local tree already avoids importing tvm.testing on the import path by using the testing.object_use_count FFI lookup directly. Keep that equivalent behavior while preserving the upstream cherry-pick point. (cherry picked from commit a979b2f)
Follow-ups to the cibuildwheel wheel-publishing flow (apache#19656): - macOS: ad-hoc re-sign the wheel's Mach-O dylibs after delocate. install_name_tool edits invalidate the arm64 code signature and dyld SIGKILLs an invalidly-signed dylib on dlopen, so `import tvm` crashed with no traceback. New ci/scripts/package/macos_repair_wheel.sh runs delocate, ad-hoc re-signs every Mach-O, and repacks so RECORD matches. - Simplify the per-platform CUDA extra-libs in the wheel CMAKE_ARGS: macOS never bundles the CUDA sidecar (drop the always-empty arg); Linux/Windows always do (pass -DTVM_PACKAGE_EXTRA_LIBS unconditionally). - Move the wheel post-install checks into tests/python/all-platform-minimal-test, gated behind TVM_WHEEL_EXPECT_LLVM / TVM_WHEEL_EXPECT_CUDA_RUNTIME so they only assert during wheel validation and skip in ordinary source-build CI; the cibuildwheel test-command is now a single pytest invocation. - Windows: collapse the two tvm_ffi DLL excludes into the delvewheel glob --exclude "*tvm_ffi*.dll" and pin delvewheel>=1.12.0 (wildcards need >=1.12.0). (cherry picked from commit da9b580)
The host/device split flow already runs device-region annotation, host/device function extraction, and device-kernel launch lowering as one consecutive pipeline. Keeping those stages exposed as separate public passes makes the API surface larger than the actual execution model and leaves the stage dependencies spread across multiple files. This change makes `tirx.transform.SplitHostDevice` the single public entry point for that flow, while preserving the existing stage order internally. Changes: - Merge the annotation, splitting, and kernel-launch lowering implementations into `src/tirx/transform/split_host_device.cc` as private sections. - Remove the old public C++ declarations, FFI registrations, and Python wrappers for `AnnotateDeviceRegions` and `LowerDeviceKernelLaunch`. - Replace pipeline call sites that previously invoked the three-stage sequence with one `SplitHostDevice()` call. - Update TIRx and S-TIR tests to exercise the consolidated pass and the reduced public API surface. (cherry picked from commit dea2bf9)
TVM can rely on tvm-ffi's JSON graph serialization helpers directly instead of routing through TVM-side `node.SaveJSON`/`node.LoadJSON` registry entries. This changes `tvm.ir` save/load to call `tvm_ffi.serialization` with `tvm_version` metadata, removes the C++ registry wrapper, and moves the disco debug object path to `ffi::ToJSONGraph`/`ffi::FromJSONGraph` plus JSON parse/stringify. The disco Python wrappers now declare Python attribute storage explicitly for `DRef` and `Session` so `DPackedFunc`/`DModule` and method caches continue to work with the current tvm-ffi object model. The socket address helper also normalizes `localhost` consistently across constructors so the disco socket debug round-trip can bind an IPv4 socket when `localhost` resolves to IPv6 first. Validated locally in an isolated worktree build with `ninja -C build tvm_compiler tvm_runtime_extra`, targeted IR/target tests, `tests/python/disco/test_session.py::test_string_obj`, import smoke, and touched-file pre-commit. (cherry picked from commit 1382707)
…pache#19660) `torch.pow` on an integer tensor returns an integer result, but the PyTorch frontend lowered it to `relax.op.power`, which fails `LegalizeOps` with `power only applies to float` (TOPI `power` / `tvm::pow` requires a floating-point input). This decomposes an integer base raised to a constant non-negative integer exponent into repeated multiplication, so the result stays integral and matches PyTorch. Float bases and non-constant or tensor exponents keep using `relax.op.power` unchanged. The ONNX frontend already uses the same decomposition (`x**3 = x*x*x`). Added structural tests covering both the FX and ExportedProgram import paths. Fixes apache#19550 (cherry picked from commit 5791239)
Replace manual version.py stamping with scikit-build-core's setuptools_scm metadata provider, so local builds no longer call version.py. The Python distribution/runtime version comes from the generated python/tvm/_version.py (libinfo.py reads it with a fallback); the C++ TVM_VERSION is injected from SKBUILD_PROJECT_VERSION_FULL with a #ifndef default in base.h for bare cmake builds. version.py is removed. The publish workflow's wheel build checks out full history (fetch-depth: 0) so setuptools_scm can derive the version, and drops the version.py stamping step. release_process.rst is updated to the tag-driven release flow. (cherry picked from commit a72d57c)
…pache#19664) The single-line 'bash -c' form is hard to read and edit; rewrite it as a readable multiline command. Functionally equivalent to the previous inline version: delocate the wheel, then ad-hoc re-sign every bundled Mach-O and repack. The re-sign step is required because delocate's edits invalidate the ad-hoc signature of tvm/lib/libtvm_runtime.dylib (it LC_LOADs the excluded `rpath/libtvm_ffi.dylib`, which ships in the separate apache-tvm-ffi package), and arm64 dyld SIGKILLs an invalidly-signed Mach-O on import. 'wheel' is installed explicitly since it is absent from cibuildwheel's repair venv, and the delocated wheel is located by glob since delocate may retag it. (cherry picked from commit 9d68a1d)
## Summary Python callers should reach the canonical tvm-ffi structural helpers directly instead of going through a TVM-side redirect layer. This makes the public tvm.ir bindings exact aliases of the tvm_ffi APIs and exposes get_first_structural_mismatch from tvm.ir. Main changes: - Import structural_equal, get_first_structural_mismatch, and structural_hash directly from tvm_ffi - Remove the pure wrappers from tvm.ir.base while keeping assert_structural_equal's TVM-specific formatting - Update mismatch tests and add identity coverage for the direct bindings (cherry picked from commit 1240649)
…blowup (apache#19670) `Analyzer::Bind` could hang indefinitely (>300s, ~200% CPU, no GPU work) while binding a small expression for one variable. The root cause is general and lives in `src/arith/int_set.cc`. Diagnosis: 100% of the time is spent in `arith::Analyzer::Bind` → `IntSetAnalyzer` → `IntervalSetEvaluator`, evaluating a **5-node** bound expression. A counter showed **>2^20 `VisitExpr` calls at recursion depth 67** with no end in sight. `IntervalSetEvaluator::VisitExpr_(VarNode)` relaxes a variable's bounds by recursively evaluating **both** the `min` and `max` sub-expressions of its mapped interval. For diamond-shaped variable dependency chains (`a → {b, c}`, `b → {d, e}`, …) the shared sub-expressions are re-expanded along every path, so cost is **O(2^depth)** in the length of the dependency chain — bounded only by `dom_map_.size()` (~67 interdependent vars in the failing case). Memoize the fully-relaxed interval **per variable** (`relax_memo_`) and break cyclic dependencies with an in-progress set (`relax_in_progress_`). A variable's relaxed interval is deterministic for a given evaluator instance (`dom_map_`/`dom_constraints_` are fixed), so memoizing collapses the diamonds to linear cost. Short chains — the common case, which never reached the old `recur_depth_ >= dom_map_.size()` cutoff — are unaffected, so the change is behavior-preserving outside the pathological case. New regression tests in `tests/python/arith/test_arith_intset.py`: - `test_relax_deep_variable_dependency_chain` — a 64-deep diamond (`O(2^64)` without the fix; verified to hang on a clean build), also asserting the relaxed result is correct (`x0 → [-n, 100+n]`). - `test_relax_cyclic_variable_dependency` — a cyclic `x↔y` dependency must terminate. - `tests/python/arith/test_arith_intset.py` — 20 passed (the deep-chain test completes instantly). - Full `tests/python/arith/` — 933 passed (1 pre-existing flaky random-seed failure in `test_arith_solve_linear_equations.py` unrelated to this change, passes on rerun). Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com> (cherry picked from commit 96b8257)
…e#19669) ## Summary `CanonicalSimplifier::Impl::VisitExpr_(LTNode)` Case 2 rewrites S + xn < 0 ⇔ S/d + xn // d < 0 where d = gcd(scales) The Case 1 derivation only works when `xn ≥ 0`. With `scale = -1` the equivalence becomes `≤` rather than `<`, and the rewrite silently strengthens the predicate by dropping the boundary `S/d == xn // d`. After CSE/inlining, a comparison such as `2*(tx%4) < 16*warp + (tx%32)//4` (where `row` and `col` are independent projections of the same lane id) reaches canonical_simplify with the divided projection on the LHS (scale = -1), and Case 2 folds it to a plain `0 < warp_id` — zeroing every thread that should have written `val` in warp 0. The same path also folds other configurations (e.g. `0 < (tx%32) - 8*warp`) all the way to `False`. The fix gates Case 2 with `extra->args[0]->scale == 1`. The original target shape (`yn % m` with positive scale and `lower_factor=1`, plus the `scale = +1 / lower_factor > 1` generalization) is unchanged; truly-always-true comparisons still fold to `True`. ## Test plan - New regression test `test_simplify_le_negative_scale_extra` in `tests/python/arith/test_arith_canonical_simplify.py` — asserts on simplified `PrimExpr`, no GPU required; pre-fix fails, post-fix passes. It also pins the buggy `scale = -1` shapes to their unsimplified form, confirms the `scale = +1` Case 2 path still optimizes, and re-asserts the truly-always-true variant still folds to `True`. - Existing `test_simplify_le` (the original Case 2 target with `scale = +1`) still passes. - `tests/python/arith/test_arith_canonical_simplify.py` — 16 passed. - Full `tests/python/arith/` — 932 passed (1 pre-existing flaky random-seed failure in `test_arith_solve_linear_equations.py` unrelated to this change, passes on rerun). (cherry picked from commit 913fc4b)
…9626) Hi Committers, This PR is trying to fix issues apache#19542. Any suggestions would be appreciated if you are available. ### Root cause: FP to INT lowering can be implementation-defined or UB for NaN/Inf and extreme floats, producing backend-dependent results versus ONNX Runtime. ### Solution: Apply a minimal, deterministic frontend sanitization for float to integer Casts: map NaN and ±Inf to 0.0 before astype. This prevents NaN/Inf from reaching backend fptosi/fptoui lowers and yields stable behavior across targets. --------- Co-authored-by: cchung100m <cchung100m@users.noreply.github.com> (cherry picked from commit 4d9d129)
## Summary - replace the block-structured TIRx exec-scope surface with scope-qualified `Tx.<scope>.<op>` namespaces and migrate call sites - split TIRx op namespaces and remove the unused dynamic generic-op fallback - add explicit CUDA launch bounds plumbing through TIRx attrs and split-host-device lowering ## Validation - `git diff --check apache/main..HEAD` - `pre-commit run --from-ref apache/main --to-ref HEAD` (cherry picked from commit 9db74c7)
…#19674) Hi Committers, This PR fixes issues apache#19543. Any suggestions would be appreciated if you are available. ### Root cause: The ONNX frontend `Sign` converter directly returned `relax.op.sign(x)`. After legalization, this maps to `topi.sign`, which is implemented via comparisons (x < 0 ? -1 : x > 0 ? 1 : 0). For `NaN`, both comparisons are false, so TVM produced 0, while ONNX Runtime preserves NaN. This created a frontend semantic mismatch for imported ONNX models. ### Solution: Apply a minimal ONNX-frontend-only fix in `onnx_frontend.py`: - For floating-point inputs, lower `Sign` as `where(isnan(x), x, sign(x))`. - Keep non-floating inputs unchanged (`sign(x)`). --------- Co-authored-by: cchung100m <cchung100m@users.noreply.github.com>
[Bump] tvm-ffi to 59da4c0 Bumps `3rdparty/tvm-ffi` from 98d0029 to 59da4c0, picking up seven commits from apache/tvm-ffi: - [FEAT] Optimize Expected<T> for minimal compiled code and efficiency (apache#599) - [FEAT] Streamline AccessPath/AccessStep print format (apache#598) - [CI] Use uv tool run for sdist build instead of pipx (apache#600) - [CORE] Make AnyView trivially copyable — match C-ABI struct passing (apache#602) - feat(python): add tensor methods to align with C++ APIs (apache#604) - [FIX] Drop test_empty_tensor_attributes (numpy dlpack stride change) (apache#607) - [TEST] Relax test_shared_dag_hash_scaling_not_exponential ratio to 4x (apache#612) The bump range is additive for almost all of the above (performance improvements to Expected<T>, new Python tensor methods, AnyView ABI alignment, CI toolchain switch). One commit required a TVM-side migration: apache/tvm-ffi#598 moved the `__ffi_repr__` hooks for `ffi::reflection::AccessPath` and `AccessStep` into tvm-ffi itself (src/ffi/extra/reflection_extra.cc, compiled into libtvm_ffi.so). TVM already registered the same hooks in src/ir/access_path_repr.cc, causing a double-registration abort at library load time. This commit removes the duplicate AccessPath/AccessStep `__ffi_repr__` registrations from access_path_repr.cc, keeping only the `node.AsRepr` global function that tvm-ffi does not provide. The format emitted by tvm-ffi's hooks is equivalent for common cases; missing-item steps now use the `[<missing:...>]` notation from tvm-ffi rather than the `...?` suffix that TVM used previously.
…ecords (apache#19673) Add support for synchronous access handles in OPFS (https://developer.mozilla.org/en-US/docs/Web/API/FileSystemFileHandle/createSyncAccessHandle). Sync mode uses FileSystemSyncAccessHandle where available. Replace optional metadata with committed OPFS records written after payload. Records store the URL, payload byte count, and content type. This allows interrupted or partial writes to be treated as cache misses. Stale OPFS directory handles are cleared on `InvalidStateError`.
This PR makes `arith::Analyzer` a first-class tvm-ffi object. The implementation splits the previous concrete `Analyzer` class into: - `AnalyzerObj`, the mutable object node that owns analyzer state, sub-analyzers, caches, and bindings - `Analyzer`, a reference-counted `ObjectRef` handle that can be passed across the tvm-ffi boundary This allows Python and C++ to share the same analyzer instance, so bindings, constraints, and cached facts can persist across FFI calls. Public APIs that accept an analyzer now use `const arith::Analyzer&`, while internal helper APIs that only borrow the object continue to use `AnalyzerObj*`. --------- Co-authored-by: Ubospica <ubospica@gmail.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
20 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Changes
tilelang_main.Validation
./maint/scripts/run_local_ci_test.sh 2>&1 | tee run.log./format.shbecause this fork does not include that script