Newer LLVM versions (observed with LLVM 20) print a scalable broadcast
store as a splat constant, e.g.
`store <vscale x 4 x float> splat (float 1.000000e+00)`, instead of the
older `shufflevector (<vscale x 4 x float> insertelement (...` form.
Accept either representation in test_scalable_broadcast so the test
passes across LLVM versions while still verifying scalable vector
codegen.
---------
Co-authored-by: tqchen <tianqi.tchen@gmail.com>
## 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.
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
This PR removes unused `from tvm import te` imports from 25 test files
across the codebase, continuing the ongoing TE → TVMScript migration
cleanup.
Changes:
- Remove unused `from tvm import te` from 24 test files in
tir-transform/, s_tir/transform/, codegen/, arith/,
all-platform-minimal-test/, and testing/
- Replace `te.var("x")` with `tvm.tir.Var("x", "int32")` in
`test_s_tir_transform_decorate_device_scope.py` (the only file where te
was actually used)
- Clean up stale `tvm.tir.ir_builder` comment references in
`test_tir_transform_convert_ssa.py`
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 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.
Shushi Hong