Repository navigation
[REFACTOR] Consume standalone TIRx as a pinned dependency - #20803
Merged
Merged
Conversation
Replace the owned TIRx tree with the Apache TIRx submodule and its recursive FFI dependency. Point native, Python, packaging, CI, and development integration at the external component while retaining TVM compiler and runtime extensions.
Clear the external FFI CMake cache in the build entry points that install it. Keep dependency-specific preparation alongside package setup.
Pin TIRx 0.1.0a2 and use its public integer-set and linear-analysis interfaces. Align build and runtime dependency requirements, refresh the matcher header reference, and simplify macOS wheel setup.
spectrometerHBH
approved these changes
Oct 20, 2026
Update the TIRx dependency and keep TOPI constant-integer extraction local after its removal from the symbolic utilities package.
This was referenced Oct 22, 2026
LngelKyo
added a commit
to LngelKyo/tvm-tirx
that referenced
this pull request
Oct 23, 2026
The CUDA codegen printed bare min()/max() for T.min/T.max on float32,
float16 and bfloat16, so a NaN operand was dropped. Emit the ordered
compare with an (a != a) test instead, bind scalar operands once, and
keep constant-operand fast paths. Tests cover 4 dtypes x 2 ops x
{scalar, vec4}, a composite expression, chained statements and the
constant-operand cases on both the nvcc and NVRTC paths.
Same change as apache/tvm#20387, re-targeted after the TIRx sources
moved to this repository in apache/tvm#20803.
LngelKyo
added a commit
to LngelKyo/tvm-tirx
that referenced
this pull request
Oct 23, 2026
The CUDA codegen printed bare min()/max() for T.min/T.max on float32,
float16 and bfloat16, so a NaN operand was dropped. Emit the ordered
compare with an (a != a) test instead, bind scalar operands once, and
keep constant-operand fast paths. Tests cover 4 dtypes x 2 ops x
{scalar, vec4}, a composite expression, chained statements and the
constant-operand cases on both the nvcc and NVRTC paths.
Same change as apache/tvm#20387, re-targeted after the TIRx sources
moved to this repository in apache/tvm#20803.
tlopex
pushed a commit
to apache/tvm-tirx
that referenced
this pull request
Oct 23, 2026
This is apache/tvm#20387 re-targeted at this repository after the TIRx sources moved here in apache/tvm#20803. The diff is the same apart from the test file path and the module paths in the tests (`tvm.compile` -> `tvm_tirx.tirx.compile`; `tvm.testing`, `tvm.runtime`, `tvm.cuda` -> the `tvm_tirx.` forms the rest of the file uses). ### Bug: on `main`, the CUDA codegen prints bare `min()`/`max()` for `T.min`/`T.max` on float32/float16/bfloat16, so a NaN operand is dropped With apache/tvm#20054's test data, the CUDA target gives 4/8 bitwise matches: the NaN lanes are discarded and the ±0 ties disagree with the C host. ### Fix Same semantics as the host-side proposal apache/tvm#20054: emit `(((a > b) || (a != a)) ? a : b)` for `max` (and the `<` form for `min`), so a NaN on either side is preserved and ties take `b`. Scalar form binds each operand once via SSA and references the bound names; the vector form expands per lane. Generated source (scalar and per-lane vector), from the branch: ``` [float32/scalar] C_ptr[((int)threadIdx.x)] = (((v_ > v__1) || (v_ != v_)) ? v_ : v__1); [float32/vec4] __1.x = ((v_.x > v__1.x) || (v_.x != v_.x)) ? v_.x : v__1.x; ``` bfloat16 constants: `PrintConst` printed a bfloat16 NaN or Inf constant as `__float2bfloat16_rn(nanf/*nan*/)` / `(inff/*inf*/)`, which does not compile (the float32 branch prints `CUDART_NAN_F` / `CUDART_INF_F`, and the float16 branch delegates to it). The bfloat16 branch now delegates to the float32 printer in the same way; finite bfloat16 constants print as before. The bfloat16 rows of the test need it; it is a separate commit. ### Tests Folded into `tests/python/backend/cuda/codegen/test_basic.py` as one parametrized test: for each dtype in {float32, float64, float16, bfloat16} × {scalar, vec4}, a kernel computes the 15-row matrix of the host-side test in #31 — min and max of (a, b), (0, a), (a, 0), (nan, a), (a, nan); `max(a, b) + 1` and `max(0, a) + 1` (which catch the ternary being parsed as `(x + cond) ? a : b`); the chained `C = 0; C = max(C, 3); C = max(C, 2)` row; `min(a, +inf)` and `max(a, −inf)` — and an int32 kernel computes min and max and must keep the base path. 10 cases, all requesting the `backend_config` fixture, so they run on both the nvcc and NVRTC paths (20 in total). For float32 the generated source is also checked: the general rows emit the `||` form, the constant NaN rows `CUDART_NAN_F`, a scalar constant rhs the reversed non-strict compare, and no bare `min()`/`max()`. The scalar path's SSA bindings must not stay alive across statements: `SSAGetID` caches by printed text within the live scope and the base-codegen `BufferStore` dispatch does not invalidate it, so `C[v] = max(C[v], A[v]); C[v] = max(C[v], B[v])` would reuse the first read of `C[v]`. Row 12 of the matrix (`C = 0; C = max(C, 3); C = max(C, 2)`) pins this: a stale cached read gives 2. Constant-operand fast paths (from the review of apache/tvm#20054): a constant NaN lhs is the result (both operands are bound before the classification, so the rhs is still evaluated); a constant NaN rhs needs only the `(a != a)` test; a constant non-NaN lhs keeps the ordered compare; a constant non-NaN rhs reverses it to the non-strict form (`<` -> `>=`, `>` -> `<=`) so ties still take `b`. The three ternary forms are parenthesized so they survive nesting in a compound expression (row 11 of the matrix); the constant-NaN-lhs form prints the constant alone. The tests are written as plain TIRx (`@T.function` with `T.constexpr` selections) for the transpiler parser. ### Verification Before the move, in apache/tvm: - RTX A6000 (sm_86) and L40 (sm_89), CUDA 13.0.88: `-k min_max` **74 passed** (nvcc + nvrtc); the whole `test_target_codegen_cuda.py` of the time **440 passed, 6 skipped**; `test_gpu_codegen_allreduce.py` **328 passed, 330 skipped**; negative control with the predicate forced to `return false`: **68 failed / 6 passed** (the 6 are the 4 int-vector and 2 chained cases; the chained ones are green by construction because the control bypasses the new path); restored: **74/74**. - apache/tvm Jenkins GPU job on the 2026-10-14 merge with main: `test_min_max` **74 passed** (37 nvcc + 37 NVRTC), 0 failed. - `ruff@0.12.3` check/format and `clang-format 20.1.8 --dry-run --Werror` clean. NaN constants and NaN results are only guaranteed to be NaN, not a particular payload (on sm_89 nvcc normalizes some forms to `0x7FFFFFFF`); the tests assert `isnan` on the NaN lanes and bitwise equality on the finite lanes. After the move: the same patch applied onto this repository's `main` (the C++ hunks apply unchanged at shifted offsets; the tests change only the module paths above). I have not re-run the GPU tests against this repository's build, so the numbers above are from apache/tvm; the Jenkins GPU pipeline here runs this file on both compiler paths. Partially addresses apache/tvm#19579 (CUDA side only). The host side (C host and LLVM codegen, plus the constant folding) is #31 in this repository; this PR does not touch those paths. The constant classification here is local to the CUDA codegen (scalar `FloatImm`); it can move to `src/codegen/min_max_utils.h` once #31 lands.
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
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.
Consume Apache TIRx as a pinned submodule at
3rdparty/tvm-tirx, replacing the owned TIRx source tree and retaining TVM's compiler and runtime extensions.Update native and Python builds, packaging, CI, development scripts, and documentation to use the standalone component and its recursive FFI dependency. Keep the parent distribution focused on TVM extensions while consuming TIRx through its installed interface.