This is an automated email from the ASF dual-hosted git repository.
tqchen pushed a change to branch tirx-standalone-root
in repository https://gitbox.apache.org/repos/asf/tvm.git
discard 40c5620ac4 [DOCS] Add standalone TIRx Python API scaffold
discard 1c423e35b7 [CI] Add standalone TIRx build, wheel, and docs pipelines
discard 6052c04140 [LINT] Add standalone TIRx lint configuration
discard 0bd0fbf0af [META] Add repository ignore rules and release keys
discard 38f24bcd89 [META] Add ASF configuration and standalone FFI submodule
mapping
discard f9efc0853f [REFACTOR] Make standalone TIRx the repository root
new 96e7e85dc7 [BUILD] Enable LLVM and automatic CUDA in standalone TIRx
(#20755)
new 69215634db [REFACTOR][TIRx] Align public APIs and package boundaries
(#20759)
new 52441cdf2b [FIX][Sym] Reject non-positive divisors in iter map
floordiv/floormod (#20756)
new 341e3fd7f5 [REFACTOR] Separate TIRx compiler and runtime deployment
APIs (#20763)
new 8a10be0a14 [TIRx] Clarify overload sets, leaf ownership, and terminal
throws (#20765)
new 80d645d824 [TIRx] Simplify Python APIs and configuration discovery
(#20766)
new 54a62ddf4b [REFACTOR] Make standalone TIRx the repository root
new 9a242ed23e [META] Add ASF configuration and standalone FFI submodule
mapping
new b2e6e56471 [META] Add repository ignore rules and release keys
new eef005bf30 [LINT] Add standalone TIRx lint configuration
new 8bfc5bd764 [CI] Add standalone TIRx build, wheel, and docs pipelines
new 6a4f354cf4 [DOCS] Add standalone TIRx Python API scaffold
This update added new revisions after undoing existing revisions.
That is to say, some revisions that were in the old version of the
branch are not in the new version. This situation occurs
when a user --force pushes a change and generates a repository
containing something like this:
* -- * -- B -- O -- O -- O (40c5620ac4)
\
N -- N -- N refs/heads/tirx-standalone-root (6a4f354cf4)
You should already have received notification emails for all of the O
revisions, and so the following emails describe only the N revisions
from the common base, B.
Any revisions marked "omit" are not gone; other references still
refer to them. Any revisions marked "discard" are gone forever.
The 12 revisions listed above as "new" are entirely new to this
repository and will be described in separate emails. The revisions
listed as "add" were already present in the repository and have only
been added to this reference.
Summary of changes:
.github/actions/build-wheel-for-publish/action.yml | 2 +-
.github/workflows/main.yml | 4 +-
.gitignore | 2 +-
CMakeLists.txt | 45 ++-
cmake/modules/CUDA.cmake | 24 +-
cmake/modules/ClangFlags.cmake | 2 +-
cmake/modules/Logging.cmake | 4 +-
docs/conf.py => cmake/tvm-tirx-config.in | 15 +-
cmake/tvm_tirx-config.cmake | 6 +-
cmake/utils/FindLLVM.cmake | 11 +-
docs/reference/python/index.rst | 9 +-
include/tvm/backend/cuda/attr.h | 20 +-
include/tvm/backend/cuda/op/iket.h | 17 +-
include/tvm/backend/cuda/op_attr_types.h | 5 +-
include/tvm/codegen/codegen.h | 6 +-
include/tvm/codegen/intrin_rule.h | 5 +-
include/tvm/ir/attrs.h | 2 +-
include/tvm/ir/expr.h | 2 +-
include/tvm/ir/location.h | 79 +---
include/tvm/ir/module.h | 26 +-
include/tvm/ir/prim/op.h | 2 +-
include/tvm/ir/transform.h | 7 +-
include/tvm/ir/with_context.h | 2 +-
include/tvm/runtime/base.h | 8 +-
include/tvm/runtime/c_backend_api.h | 9 +-
include/tvm/runtime/device_api.h | 8 +-
include/tvm/runtime/memory_scope.h | 17 +-
include/tvm/runtime/tensor.h | 106 +-----
include/tvm/runtime/timer.h | 120 +------
include/tvm/script/ir_builder/base.h | 4 +
include/tvm/script/ir_builder/frame.h | 21 +-
include/tvm/script/ir_builder/ir.h | 4 +
include/tvm/script/printer/doc.h | 4 +
include/tvm/script/printer/doc_translator.h | 5 +
include/tvm/support/base64.h | 302 ----------------
include/tvm/support/bytes_io.h | 2 +-
include/tvm/support/cuda/nvtx.h | 71 ----
include/tvm/support/env.h | 2 +-
include/tvm/support/ordered_map.h | 2 +-
include/tvm/support/str_escape.h | 123 -------
include/tvm/sym/bound.h | 25 +-
include/tvm/sym/int_set.h | 31 +-
include/tvm/target/target.h | 6 +-
include/tvm/target/target_kind.h | 11 +-
include/tvm/tirx/analysis.h | 21 +-
include/tvm/tirx/function.h | 20 +-
include/tvm/tirx/index_map.h | 6 +-
include/tvm/tirx/layout.h | 3 -
include/tvm/tirx/op/memory.h | 10 +-
include/tvm/tirx/op/region.h | 3 +-
include/tvm/tirx/op_attr_types.h | 10 +-
include/tvm/tirx/script/ir_builder/frame.h | 20 +-
include/tvm/tirx/script/ir_builder/ir.h | 89 +----
include/tvm/tirx/stmt.h | 10 +-
include/tvm/tirx/stmt_functor.h | 2 +-
include/tvm/tirx/tile/context.h | 5 +-
include/tvm/tirx/tile/dispatch.h | 5 +-
include/tvm/tirx/tile/op.h | 5 +-
include/tvm/tirx/tile/thread_scope.h | 5 +-
include/tvm/tirx/transform.h | 6 +-
pyproject.toml | 8 +-
python/tvm_tirx/__init__.py | 31 +-
python/tvm_tirx/backend/cuda/codegen/registry.py | 26 +-
python/tvm_tirx/backend/cuda/codegen/schema.py | 6 +-
python/tvm_tirx/backend/cuda/cpp/asm.py | 40 +--
python/tvm_tirx/backend/cuda/cpp/builtins.py | 86 ++---
python/tvm_tirx/backend/cuda/cpp/descriptors.py | 12 +-
python/tvm_tirx/backend/cuda/cpp/instrument.py | 14 +-
python/tvm_tirx/backend/cuda/cpp/nvshmem.py | 6 +-
python/tvm_tirx/backend/cuda/iket.py | 2 +-
python/tvm_tirx/backend/cuda/lang/pipeline.py | 2 +-
python/tvm_tirx/backend/cuda/launch/_impl.py | 2 +-
python/tvm_tirx/backend/cuda/op.py | 112 +++---
python/tvm_tirx/backend/cuda/ptx/engine.py | 20 +-
python/tvm_tirx/backend/cuda/ptx/gen_stubs.py | 6 +-
python/tvm_tirx/backend/cuda/script.py | 48 +--
.../backend/cuda/tile/patterns/async_copy.py | 10 +-
.../tvm_tirx/backend/cuda/tile/patterns/common.py | 5 +-
.../backend/cuda/tile/patterns/descriptors.py | 9 +-
.../backend/cuda/tile/patterns/tcgen05_mma.py | 8 +-
python/tvm_tirx/backend/cuda/tile/patterns/tma.py | 14 +-
python/tvm_tirx/backend/metal/op.py | 15 +-
python/tvm_tirx/backend/metal/target_tags.py | 4 +-
python/tvm_tirx/backend/trn/tile/binary/default.py | 3 +-
python/tvm_tirx/backend/trn/tile/binary/utils.py | 3 +-
.../backend/trn/tile/instruction_generator.py | 4 +-
python/tvm_tirx/backend/trn/tile/private_alloc.py | 4 +-
python/tvm_tirx/backend/trn/tile/select/default.py | 4 +-
python/tvm_tirx/backend/trn/tile/unary/default.py | 3 +-
python/tvm_tirx/backend/trn/tile/unary/utils.py | 3 +-
.../backend/trn/transform/naive_allocator.py | 3 +-
python/tvm_tirx/base.py | 8 +-
python/tvm_tirx/codegen/__init__.py | 68 ++++
.../{runtime/module.py => codegen/executable.py} | 351 +++++-------------
python/tvm_tirx/{codegen.py => codegen/llvm.py} | 72 +---
python/tvm_tirx/config.py | 211 ++++++-----
python/tvm_tirx/error.py | 126 -------
python/tvm_tirx/ir/__init__.py | 3 +
python/tvm_tirx/ir/_overload_prim_expr.py | 2 +-
python/tvm_tirx/ir/attrs.py | 38 +-
python/tvm_tirx/ir/base.py | 4 -
python/tvm_tirx/ir/expr.py | 114 +++++-
python/tvm_tirx/ir/function.py | 7 +-
python/tvm_tirx/ir/json_compact.py | 64 ----
python/tvm_tirx/ir/location.py | 8 +-
python/tvm_tirx/ir/module.py | 38 +-
python/tvm_tirx/ir/op.py | 234 ++++++------
python/tvm_tirx/ir/prim/__init__.py | 6 +-
python/tvm_tirx/ir/prim/expr.py | 171 +--------
python/tvm_tirx/ir/prim/op.py | 96 ++---
python/tvm_tirx/ir/prim/vector_expr.py | 109 ++++++
python/tvm_tirx/ir/stmt.py | 12 +-
python/tvm_tirx/ir/type.py | 7 +-
.../ir/{supply.py => unique_name_supply.py} | 0
python/tvm_tirx/libinfo.py | 149 +++++++-
python/tvm_tirx/runtime/__init__.py | 43 ++-
python/tvm_tirx/runtime/_tensor.py | 10 +-
python/tvm_tirx/runtime/device.py | 19 +-
python/tvm_tirx/runtime/executable.py | 175 ---------
python/tvm_tirx/script/ir_builder/_ffi_api.py | 22 --
python/tvm_tirx/script/ir_builder/frame.py | 14 +-
python/tvm_tirx/script/ir_builder/stmt.py | 3 +-
python/tvm_tirx/script/printer/doc.py | 2 +-
python/tvm_tirx/script/tirx.pyi | 2 +-
python/tvm_tirx/support/cc.py | 6 +-
python/tvm_tirx/support/clang.py | 2 +-
python/tvm_tirx/support/ndk.py | 4 +-
python/tvm_tirx/support/rocm.py | 2 +-
python/tvm_tirx/sym/analyzer.py | 2 +-
python/tvm_tirx/sym/bound.py | 2 +-
python/tvm_tirx/sym/int_set.py | 10 +-
python/tvm_tirx/sym/iter_affine_map.py | 18 +-
python/tvm_tirx/sym/utils.py | 2 +-
python/tvm_tirx/target/intrin.py | 81 -----
python/tvm_tirx/target/target.py | 14 +-
python/tvm_tirx/testing/env.py | 8 +-
python/tvm_tirx/tirx/__init__.py | 10 +-
python/tvm_tirx/tirx/_compile.py | 4 +-
python/tvm_tirx/tirx/_tensor_view.py | 8 +-
python/tvm_tirx/tirx/analysis/analysis.py | 2 +-
python/tvm_tirx/tirx/backend/__init__.py | 17 -
python/tvm_tirx/tirx/expr.py | 39 --
python/tvm_tirx/tirx/function.py | 297 +--------------
python/tvm_tirx/tirx/{function.py => index_map.py} | 169 +--------
python/tvm_tirx/tirx/layout.py | 5 +-
python/tvm_tirx/tirx/op/__init__.py | 249 +++++--------
python/tvm_tirx/tirx/script/ir_builder/frame.py | 7 -
python/tvm_tirx/tirx/script/ir_builder/ir.py | 19 +-
python/tvm_tirx/tirx/script/ir_builder/op.py | 53 ++-
.../tirx/script/ir_builder/parser_protocol.py | 67 ++--
python/tvm_tirx/tirx/script/ir_builder/triton.py | 3 +-
python/tvm_tirx/tirx/tile/instruction.py | 6 +-
src/backend/cuda/codegen/codegen_cuda_host.cc | 2 +
src/backend/cuda/runtime/cuda_device_api.cc | 2 +-
src/backend/cuda/transforms/bind_backend_config.cc | 6 +-
src/backend/cuda/transforms/lower_iket.cc | 7 +-
src/backend/metal/runtime/metal_device_api.mm | 2 +-
src/backend/opencl/runtime/opencl_common.h | 2 +-
src/backend/rocm/runtime/rocm_device_api.cc | 2 +-
src/backend/vulkan/runtime/vulkan_timer.h | 2 +-
src/codegen/c/codegen_c.h | 1 +
src/codegen/c/codegen_c_host.cc | 13 +-
src/codegen/llvm/codegen_cpu.cc | 3 +-
src/codegen/llvm/codegen_llvm.h | 1 +
src/codegen/llvm/codegen_x86_64.cc | 2 +
src/ir/attrs.cc | 5 -
src/ir/env_func.cc | 13 +-
src/ir/expr.cc | 30 +-
src/ir/function.cc | 1 -
src/ir/location.cc | 70 ----
src/ir/module.cc | 49 ++-
src/ir/op.cc | 12 +-
src/ir/prim/math.cc | 1 +
src/ir/prim/op.cc | 15 +-
src/ir/type.cc | 1 -
src/runtime/device_api.cc | 56 +--
src/runtime/tensor.cc | 116 ++++--
src/runtime/timer.cc | 276 --------------
src/script/ir_builder/frame.cc | 4 +-
src/script/ir_builder/ir.cc | 28 +-
src/script/printer/doc_printer.cc | 139 +++-----
src/script/printer/doc_translator.cc | 17 -
src/script/printer/printer.cc | 2 +-
src/sym/analyzer.cc | 3 +-
src/sym/bound_deducer.cc | 34 +-
src/sym/const_int_bound.cc | 2 +-
src/sym/int_set.cc | 34 +-
src/sym/iter_affine_map.cc | 27 +-
src/sym/modular_set.cc | 3 +-
src/sym/z3_prover.cc | 2 +-
src/target/target.cc | 10 -
src/tirx/analysis/check_well_formed.cc | 1 +
src/tirx/analysis/verify_memory.cc | 2 +-
src/tirx/ir/data_type_rewriter.h | 6 +-
src/tirx/ir/index_map.cc | 4 +-
src/tirx/ir/ir_mutator_with_analyzer.cc | 2 +-
src/tirx/op/region.cc | 13 +-
src/tirx/script/ir_builder/ir.cc | 138 +------
src/tirx/tile/lower.cc | 4 +-
src/tirx/transform/dtype_conversion.h | 6 +-
src/tirx/transform/force_narrow_index_to_i32.cc | 2 +-
src/tirx/transform/ir_utils.h | 6 +-
src/tirx/transform/narrow_datatype.cc | 2 +
src/tirx/transform/remove_no_op.cc | 9 +-
src/tirx/transform/remove_no_op.h | 6 +-
src/tirx/transform/split_host_device.cc | 32 +-
src/tirx/transform/tvm_ffi_binder.h | 6 +-
src/tirx/transform/vectorize_loop.cc | 4 +
tests/cpp/CMakeLists.txt | 2 +-
tests/cpp/runtime/backend_workspace_test.cc | 54 +++
tests/cpp/runtime/tensor_io_test.cc | 105 ++++++
tests/cpp/script/string_escape_test.cc | 65 ++++
tests/cpp/sym/domain_map_test.cc | 66 ++++
tests/python/backend/cuda/codegen/test_basic.py | 2 +-
tests/python/backend/cuda/codegen/test_fastmath.py | 2 +-
tests/python/backend/metal/codegen/test_basic.py | 12 +-
tests/python/backend/vulkan/codegen/test_basic.py | 6 +-
tests/python/backend/webgpu/codegen/test_basic.py | 20 +-
tests/python/codegen/c/test_basic.py | 8 +-
tests/python/codegen/llvm/test_aarch64.py | 36 +-
tests/python/codegen/llvm/test_basic.py | 22 +-
tests/python/codegen/llvm/test_riscv.py | 2 +-
tests/python/codegen/llvm/test_scalable_vector.py | 8 +-
tests/python/codegen/llvm/test_x86.py | 2 +-
tests/python/codegen/test_executable.py | 145 ++++++++
tests/python/conftest.py | 78 ++++
tests/python/ir/test_attrs.py | 29 ++
tests/python/ir/test_expr_layout.py | 80 +++++
tests/python/ir/test_json.py | 19 +-
tests/python/ir/test_module.py | 2 +-
tests/python/runtime/test_module.py | 152 +-------
tests/python/runtime/test_tirx_config.py | 46 ++-
.../test_validate_wheel_library.py} | 45 ++-
tests/python/script/test_basic_usage.py | 4 +-
tests/python/script/test_error_handling.py | 3 +-
tests/python/script/test_ir_builder.py | 9 +-
tests/python/script/test_printer_doc.py | 4 +-
tests/python/script/test_source_locations.py | 12 +-
tests/python/support/test_config.py | 131 ++++++-
tests/python/sym/test_analyzer_object.py | 25 --
tests/python/sym/test_const_int_bound.py | 6 +-
tests/python/sym/test_intset.py | 8 +-
tests/python/sym/test_iter_affine_map.py | 41 ++-
tests/python/sym/test_rewrite_simplify.py | 397 +++++++++++++++------
tests/python/sym/test_z3.py | 6 +-
tests/python/target/test_llvm_features_info.py | 34 +-
tests/python/target/test_riscv_features.py | 6 +-
tests/python/target/test_x86_features.py | 8 +-
tests/python/testing/test_env.py | 332 -----------------
.../tirx/analysis/test_check_well_defined.py | 26 +-
tests/python/tirx/base/test_imm_values.py | 24 +-
tests/python/tirx/base/test_index_map.py | 4 +-
tests/python/tirx/base/test_nodes.py | 48 +--
tests/python/tirx/base/test_op_types.py | 111 ++++--
tests/python/tirx/base/test_scalable_datatype.py | 10 +-
tests/python/tirx/codegen/test_cuda.py | 20 +-
tests/python/tirx/codegen/test_cuda_wait_until.py | 4 +-
tests/python/tirx/codegen/test_dsmem.py | 8 +-
tests/python/tirx/codegen/test_hopper.py | 2 +-
tests/python/tirx/codegen/test_ptx_addr.py | 8 +-
tests/python/tirx/codegen/test_ptx_dialect.py | 10 +-
.../codegen/test_tcgen05_descriptor_encoders.py | 4 +-
tests/python/tirx/iket/test_profiler.py | 9 +-
tests/python/tirx/script/test_basic_usage.py | 16 +-
tests/python/tirx/script/test_gpu.py | 14 +-
tests/python/tirx/script/test_source_locations.py | 22 +-
tests/python/tirx/test_cuda_backend_config.py | 2 +-
tests/python/tirx/test_op.py | 2 +-
tests/python/tirx/test_op_namespace_cleanup.py | 11 +-
.../python/tirx/tile/cuda/test_add_sub_mul_div.py | 6 +-
tests/python/tirx/tile/cuda/test_cp_async_bulk.py | 4 +-
.../tirx/tile/cuda/test_cp_async_bulk_tensor.py | 6 +-
tests/python/tirx/tile/cuda/test_cvt.py | 6 +-
tests/python/tirx/tile/cuda/test_ewise.py | 22 +-
tests/python/tirx/tile/cuda/test_fma.py | 4 +-
.../python/tirx/tile/cuda/test_ld_st_registers.py | 26 +-
tests/python/tirx/tile/cuda/test_mma_sync.py | 4 +-
tests/python/tirx/tile/cuda/test_mov.py | 2 +-
tests/python/tirx/tile/cuda/test_tcgen05_cp.py | 6 +-
tests/python/tirx/tile/cuda/test_tcgen05_ld_st.py | 6 +-
tests/python/tirx/tile/test_dispatcher.py | 5 +-
tests/python/tirx/transform/test_flatten_tensor.py | 4 +-
.../transform/test_force_narrow_index_to_i32.py | 12 +-
tests/python/tirx/transform/test_hoist_if.py | 2 +-
tests/python/tirx/transform/test_lower_tirx.py | 23 +-
.../python/tirx/transform/test_make_packed_api.py | 6 +-
.../python/tirx/transform/test_narrow_datatype.py | 6 +-
tests/python/tirx/transform/test_simplify.py | 5 +-
.../tirx/transform/test_split_host_device.py | 2 +-
tests/python/tirx/transform/test_unroll.py | 2 +-
tests/python/tirx/transform/test_vectorize.py | 16 +-
291 files changed, 3563 insertions(+), 5444 deletions(-)
copy docs/conf.py => cmake/tvm-tirx-config.in (71%)
delete mode 100644 include/tvm/support/base64.h
delete mode 100644 include/tvm/support/cuda/nvtx.h
delete mode 100644 include/tvm/support/str_escape.h
create mode 100644 python/tvm_tirx/codegen/__init__.py
rename python/tvm_tirx/{runtime/module.py => codegen/executable.py} (52%)
rename python/tvm_tirx/{codegen.py => codegen/llvm.py} (75%)
delete mode 100644 python/tvm_tirx/error.py
create mode 100644 python/tvm_tirx/ir/prim/vector_expr.py
rename python/tvm_tirx/ir/{supply.py => unique_name_supply.py} (100%)
delete mode 100644 python/tvm_tirx/runtime/executable.py
delete mode 100644 python/tvm_tirx/target/intrin.py
delete mode 100644 python/tvm_tirx/tirx/backend/__init__.py
copy python/tvm_tirx/tirx/{function.py => index_map.py} (69%)
delete mode 100644 src/runtime/timer.cc
create mode 100644 tests/cpp/runtime/backend_workspace_test.cc
create mode 100644 tests/cpp/runtime/tensor_io_test.cc
create mode 100644 tests/cpp/script/string_escape_test.cc
create mode 100644 tests/cpp/sym/domain_map_test.cc
create mode 100644 tests/python/codegen/test_executable.py
create mode 100644 tests/python/conftest.py
create mode 100644 tests/python/ir/test_expr_layout.py
rename
tests/python/{all-platform-minimal-test/test_validate_runtime_library.py =>
runtime/test_validate_wheel_library.py} (63%)
delete mode 100644 tests/python/testing/test_env.py
