This is an automated email from the ASF dual-hosted git repository.
github-actions[bot] pushed a change to branch nightly
in repository https://gitbox.apache.org/repos/asf/tvm.git
from 093d785e22 [Relax][PyTorch] Add atan2 converter (#19850)
add 1bb5cf6102 [REFACTOR][IR] Unify StructInfo and Type (#19853)
add 15f62fe020 [Docs] Add TIRx documentation section (#19855)
add b64cb5b2c3 [Metal] Enable Metal 4 shader compilation (#19595)
add a9ce41e5a2 [Relax][Frontend][TFLite] Add explicit operator marker
handling (#19824)
No new revisions were added by this update.
Summary of changes:
docs/arch/fusion.rst | 6 +-
docs/arch/tvmscript.rst | 24 +-
docs/conf.py | 5 +
docs/deep_dive/relax/dpl.rst | 14 +-
docs/deep_dive/relax/learning.rst | 22 +-
docs/deep_dive/relax/tutorials/relax_creation.py | 4 +-
.../tutorials/mix_python_and_tvm_with_pymodule.py | 18 +-
docs/index.rst | 12 +
docs/reference/api/python/index.rst | 10 -
docs/reference/security.rst | 6 +-
.../api/python/tirx => tirx/api}/analysis.rst | 0
.../relax/frontend.rst => tirx/api/backend.rst} | 54 +-
.../api/python/driver.rst => tirx/api/index.rst} | 16 +-
.../api/python/tirx => tirx/api}/stmt_functor.rst | 0
.../api/python/tirx => tirx/api}/tirx.rst | 2 +-
.../api/python/tirx => tirx/api}/transform.rst | 2 +-
.../tirx/backend.rst => tirx/arch/index.rst} | 15 +-
docs/tirx/arch/lowering_pipeline.rst | 203 ++
docs/tirx/install.rst | 100 +
docs/tirx/layout.rst | 383 ++++
docs/tirx/native_basics.rst | 69 +
docs/tirx/native_basics/cuda/buffers.rst | 477 +++++
docs/tirx/native_basics/cuda/compiling.rst | 112 +
docs/tirx/native_basics/cuda/control_flow.rst | 100 +
docs/tirx/native_basics/cuda/data_types.rst | 116 +
docs/tirx/native_basics/cuda/first_kernel.rst | 93 +
docs/tirx/native_basics/cuda/functions.rst | 351 +++
docs/tirx/native_basics/cuda/parser_utils.rst | 80 +
docs/tirx/native_basics/cuda/threads_sync.rst | 95 +
docs/tirx/overview.rst | 196 ++
docs/tirx/tile_primitives.rst | 297 +++
docs/tirx/tile_primitives/copy.rst | 57 +
docs/tirx/tile_primitives/copy/fallback.rst | 158 ++
docs/tirx/tile_primitives/copy/gmem_smem.rst | 209 ++
docs/tirx/tile_primitives/copy/ldstmatrix.rst | 218 ++
docs/tirx/tile_primitives/copy/reg.rst | 206 ++
docs/tirx/tile_primitives/copy_async.rst | 63 +
docs/tirx/tile_primitives/copy_async/dsmem.rst | 190 ++
docs/tirx/tile_primitives/copy_async/ldgsts.rst | 176 ++
.../tirx/tile_primitives/copy_async/tcgen05_cp.rst | 183 ++
.../tile_primitives/copy_async/tcgen05_ldst.rst | 172 ++
docs/tirx/tile_primitives/copy_async/tma.rst | 200 ++
docs/tirx/tile_primitives/elementwise.rst | 48 +
docs/tirx/tile_primitives/elementwise/reg.rst | 133 ++
docs/tirx/tile_primitives/elementwise/smem.rst | 136 ++
docs/tirx/tile_primitives/gemm.rst | 176 ++
docs/tirx/tile_primitives/gemm_async.rst | 175 ++
docs/tirx/tile_primitives/permute_layout.rst | 163 ++
.../tile_primitives/reduction.rst} | 41 +-
docs/tirx/tile_primitives/reduction/local.rst | 136 ++
docs/tirx/tile_primitives/reduction/shared.rst | 136 ++
.../tile_primitives/reduction/sm100_packed.rst | 147 ++
include/tvm/ir/expr.h | 72 +-
include/tvm/ir/op.h | 3 -
include/tvm/ir/type.h | 62 +-
include/tvm/relax/analysis.h | 182 +-
include/tvm/relax/attrs/distributed.h | 2 +-
include/tvm/relax/binding_rewrite.h | 4 +-
include/tvm/relax/block_builder.h | 6 +-
include/tvm/relax/dataflow_pattern.h | 35 +-
include/tvm/relax/dataflow_pattern_functor.h | 7 +-
include/tvm/relax/distributed/axis_group_graph.h | 2 +-
.../relax/distributed/{struct_info.h => type.h} | 56 +-
include/tvm/relax/expr.h | 155 +-
include/tvm/relax/expr_functor.h | 102 +-
include/tvm/relax/nested_msg.h | 59 +-
include/tvm/relax/op_attr_types.h | 7 +-
include/tvm/relax/script/builder/frame.h | 14 +-
include/tvm/relax/script/builder/ir.h | 24 +-
include/tvm/relax/struct_info.h | 424 ----
include/tvm/relax/struct_info_functor.h | 160 --
include/tvm/relax/transform.h | 2 +-
include/tvm/relax/type.h | 405 +++-
include/tvm/{ir => relax}/type_functor.h | 84 +-
include/tvm/relax/utils.h | 15 +-
include/tvm/script/printer/config.h | 4 +-
python/tvm/backend/cuda/op.py | 2 +-
python/tvm/contrib/cutlass/build.py | 32 +-
python/tvm/contrib/hexagon/generate_take_op.py | 6 +-
python/tvm/contrib/hexagon/hexagon_unary_ops.py | 14 +-
python/tvm/ir/expr.py | 13 +-
python/tvm/ir/global_info.py | 2 +-
python/tvm/ir/type.py | 4 +-
python/tvm/relax/__init__.py | 16 +-
python/tvm/relax/analysis/__init__.py | 10 +-
python/tvm/relax/analysis/analysis.py | 138 +-
python/tvm/relax/backend/adreno/clml.py | 40 +-
.../relax/backend/contrib/example_npu/patterns.py | 2 +-
python/tvm/relax/backend/cuda/cublas.py | 14 +-
python/tvm/relax/backend/cuda/cudnn.py | 14 +-
python/tvm/relax/backend/cuda/cutlass.py | 46 +-
python/tvm/relax/backend/dispatch_sampling.py | 6 +-
python/tvm/relax/backend/dispatch_sort_scan.py | 26 +-
python/tvm/relax/backend/metal/coreml.py | 18 +-
python/tvm/relax/backend/patterns.py | 26 +-
python/tvm/relax/backend/rocm/hipblas.py | 12 +-
python/tvm/relax/backend/utils.py | 22 +-
python/tvm/relax/base_py_module.py | 32 +-
python/tvm/relax/block_builder.py | 36 +-
python/tvm/relax/distributed/__init__.py | 2 +-
.../relax/distributed/{struct_info.py => type.py} | 22 +-
python/tvm/relax/dpl/pattern.py | 22 +-
python/tvm/relax/dpl/rewrite.py | 2 +-
python/tvm/relax/expr.py | 178 +-
python/tvm/relax/expr_functor.py | 10 +-
python/tvm/relax/frontend/common.py | 4 +-
python/tvm/relax/frontend/nn/core.py | 66 +-
python/tvm/relax/frontend/nn/exporter.py | 14 +-
python/tvm/relax/frontend/nn/extern.py | 4 +-
python/tvm/relax/frontend/nn/llm/kv_cache.py | 50 +-
python/tvm/relax/frontend/nn/modules.py | 10 +-
python/tvm/relax/frontend/nn/op.py | 24 +-
python/tvm/relax/frontend/nn/subroutine.py | 30 +-
python/tvm/relax/frontend/onnx/onnx_frontend.py | 358 ++--
.../frontend/stablehlo/stablehlo_translator.py | 10 +-
.../tvm/relax/frontend/tflite/tflite_frontend.py | 74 +-
.../frontend/torch/base_fx_graph_translator.py | 178 +-
.../frontend/torch/exported_program_translator.py | 58 +-
python/tvm/relax/frontend/torch/fx_translator.py | 76 +-
python/tvm/relax/ir/instrument.py | 10 +-
python/tvm/relax/op/_op_gradient.py | 36 +-
python/tvm/relax/op/base.py | 179 +-
python/tvm/relax/op/binary.py | 6 +-
python/tvm/relax/op/distributed/distributed.py | 21 +-
python/tvm/relax/op/grad/grad.py | 8 +-
python/tvm/relax/op/manipulate.py | 2 +-
python/tvm/relax/op/mask.py | 2 +-
.../tvm/relax/op/vision/multibox_transform_loc.py | 4 +-
python/tvm/relax/relax_to_pyfunc_converter.py | 12 +-
python/tvm/relax/script/builder/distributed/ir.py | 30 +-
python/tvm/relax/script/builder/ir.py | 150 +-
python/tvm/relax/script/parser/dist.py | 22 +-
python/tvm/relax/script/parser/entry.py | 146 +-
python/tvm/relax/script/parser/parser.py | 78 +-
python/tvm/relax/testing/ast_printer.py | 94 +-
python/tvm/relax/testing/attention.py | 2 +-
python/tvm/relax/testing/nn.py | 6 +-
python/tvm/relax/training/loss.py | 64 +-
python/tvm/relax/training/optimizer.py | 52 +-
python/tvm/relax/training/setup_trainer.py | 6 +-
python/tvm/relax/training/trainer.py | 10 +-
python/tvm/relax/transform/__init__.py | 2 +-
.../fold_batch_norm_to_conv2d_for_inference.py | 6 +-
.../tvm/relax/transform/fuse_transpose_matmul.py | 12 +-
.../tvm/relax/transform/lazy_transform_params.py | 40 +-
.../transform/legalize_ops/adreno/convolution.py | 2 +-
python/tvm/relax/transform/legalize_ops/ccl.py | 38 +-
python/tvm/relax/transform/legalize_ops/common.py | 4 +-
python/tvm/relax/transform/legalize_ops/create.py | 18 +-
.../relax/transform/legalize_ops/distributed.py | 10 +-
python/tvm/relax/transform/legalize_ops/index.py | 12 +-
.../relax/transform/legalize_ops/linear_algebra.py | 22 +-
.../tvm/relax/transform/legalize_ops/manipulate.py | 32 +-
python/tvm/relax/transform/legalize_ops/nn.py | 22 +-
python/tvm/relax/transform/legalize_ops/search.py | 4 +-
.../relax/transform/legalize_ops/statistical.py | 8 +-
python/tvm/relax/transform/legalize_ops/vision.py | 2 +-
.../relax/transform/lower_gpu_ipc_alloc_storage.py | 6 +-
.../relax/transform/optimize_layout_transform.py | 4 +-
.../relax/transform/remove_redundant_reshape.py | 4 +-
python/tvm/relax/transform/transform.py | 24 +-
python/tvm/relax/ty.py | 45 +-
python/tvm/relax/{struct_info.py => type.py} | 112 +-
python/tvm/relax/utils.py | 44 +-
python/tvm/runtime/script_printer.py | 16 +-
python/tvm/s_tir/dlight/benchmark/bench.py | 8 +-
python/tvm/s_tir/dlight/benchmark/extract.py | 22 +-
python/tvm/s_tir/dlight/benchmark/utils.py | 8 +-
python/tvm/tirx/op.py | 8 +-
src/backend/metal/runtime/metal_module.mm | 27 +-
src/ir/module.cc | 1 -
src/ir/type.cc | 3 +-
src/ir/type_functor.cc | 93 -
src/relax/analysis/analysis.cc | 4 +-
src/relax/analysis/computable_at_compile_time.cc | 2 +-
.../{struct_info_analysis.cc => type_analysis.cc} | 557 +++--
src/relax/analysis/well_formed.cc | 163 +-
.../backend/adreno/annotate_custom_storage.cc | 205 +-
.../backend/adreno/fold_vdevice_scope_change.cc | 19 +-
src/relax/backend/contrib/codegen_c/codegen_c.h | 24 +-
.../backend/contrib/codegen_json/codegen_json.h | 33 +-
src/relax/backend/contrib/cublas/codegen.cc | 6 +-
src/relax/backend/contrib/cudnn/codegen.cc | 9 +-
src/relax/backend/contrib/cutlass/codegen.cc | 18 +-
src/relax/backend/contrib/tensorrt/codegen.cc | 6 +-
src/relax/backend/vm/codegen_vm.cc | 8 +-
src/relax/backend/vm/codegen_vm_tir.cc | 2 +-
src/relax/backend/vm/lower_runtime_builtin.cc | 36 +-
src/relax/backend/vm/vm_shape_lower.cc | 167 +-
src/relax/distributed/axis_group_graph.cc | 72 +-
.../distributed/transform/legalize_redistribute.cc | 10 +-
src/relax/distributed/transform/lower_distir.cc | 95 +-
.../transform/lower_global_view_to_local_view.cc | 28 +-
.../distributed/transform/propagate_sharding.cc | 155 +-
src/relax/distributed/transform/utils.cc | 30 +-
src/relax/distributed/transform/utils.h | 8 +-
src/relax/distributed/{struct_info.cc => type.cc} | 29 +-
src/relax/ir/block_builder.cc | 194 +-
src/relax/ir/dataflow_block_rewriter.cc | 2 +-
src/relax/ir/dataflow_expr_rewriter.cc | 28 +-
src/relax/ir/dataflow_matcher.cc | 22 +-
src/relax/ir/dataflow_matcher.h | 4 +-
src/relax/ir/dataflow_pattern.cc | 27 +-
src/relax/ir/dataflow_pattern_functor.cc | 4 +-
src/relax/ir/dependent_type.cc | 219 ++
src/relax/ir/emit_te.cc | 13 +-
src/relax/ir/emit_te.h | 2 +-
src/relax/ir/expr.cc | 282 ++-
src/relax/ir/expr_functor.cc | 215 +-
src/relax/ir/py_expr_functor.cc | 8 +-
src/relax/ir/struct_info.cc | 244 ---
src/relax/ir/struct_info_functor.cc | 156 --
src/relax/ir/transform.cc | 14 +-
src/relax/ir/type.cc | 54 +-
src/relax/ir/type_functor.cc | 154 ++
src/relax/op/ccl/ccl.cc | 43 +-
src/relax/op/distributed/binary.cc | 60 +-
src/relax/op/distributed/binary.h | 47 +-
src/relax/op/distributed/ccl.cc | 20 +-
src/relax/op/distributed/distributed.cc | 88 +-
src/relax/op/distributed/linear_algebra.cc | 33 +-
src/relax/op/distributed/linear_algebra.h | 4 +-
src/relax/op/distributed/manipulate.cc | 76 +-
src/relax/op/distributed/manipulate.h | 6 +-
src/relax/op/distributed/nn.cc | 30 +-
src/relax/op/distributed/nn.h | 4 +-
src/relax/op/distributed/op.cc | 18 +-
src/relax/op/distributed/statistical.cc | 39 +-
src/relax/op/distributed/statistical.h | 10 +-
src/relax/op/distributed/unary.cc | 66 +-
src/relax/op/distributed/unary.h | 50 +-
src/relax/op/distributed/utils.cc | 71 +-
src/relax/op/distributed/utils.h | 20 +-
src/relax/op/image/resize.cc | 104 +-
src/relax/op/memory/view.cc | 125 +-
src/relax/op/nn/attention.cc | 44 +-
src/relax/op/nn/convolution.cc | 174 +-
src/relax/op/nn/nn.cc | 427 ++--
src/relax/op/nn/nn.h | 4 +-
src/relax/op/nn/pooling.cc | 138 +-
src/relax/op/op.cc | 629 +++---
src/relax/op/op_common.cc | 49 +-
src/relax/op/op_common.h | 209 +-
src/relax/op/tensor/binary.cc | 96 +-
src/relax/op/tensor/binary.h | 12 +-
src/relax/op/tensor/create.cc | 122 +-
src/relax/op/tensor/datatype.cc | 26 +-
src/relax/op/tensor/grad.cc | 42 +-
src/relax/op/tensor/index.cc | 235 +--
src/relax/op/tensor/inspect.cc | 166 +-
src/relax/op/tensor/inspect.h | 32 +-
src/relax/op/tensor/linear_algebra.cc | 127 +-
src/relax/op/tensor/manipulate.cc | 1153 +++++-----
src/relax/op/tensor/qdq.cc | 132 +-
src/relax/op/tensor/sampling.cc | 54 +-
src/relax/op/tensor/search.cc | 115 +-
src/relax/op/tensor/set.cc | 69 +-
src/relax/op/tensor/sorting.cc | 53 +-
src/relax/op/tensor/statistical.cc | 99 +-
src/relax/op/tensor/statistical.h | 2 +-
src/relax/op/tensor/ternary.cc | 32 +-
src/relax/op/tensor/unary.cc | 8 +-
src/relax/op/tensor/unary.h | 10 +-
src/relax/op/vision/multibox_transform_loc.cc | 64 +-
src/relax/op/vision/nms.cc | 167 +-
src/relax/op/vision/roi_align.cc | 30 +-
src/relax/op/vision/roi_pool.cc | 30 +-
src/relax/script/builder/distributed.cc | 29 +-
src/relax/script/builder/frame.cc | 4 +-
src/relax/script/builder/ir.cc | 43 +-
src/relax/script/builder/utils.h | 6 +-
src/relax/script/printer/binding.cc | 18 +-
src/relax/script/printer/call.cc | 36 +-
.../printer/{struct_info.cc => dependent_type.cc} | 72 +-
src/relax/script/printer/distributed.cc | 24 +-
src/relax/script/printer/expr.cc | 8 +-
src/relax/script/printer/function.cc | 20 +-
src/relax/script/printer/type.cc | 25 -
src/relax/script/printer/utils.h | 41 +-
src/relax/training/utils.cc | 18 +-
src/relax/transform/adjust_matmul_order.cc | 8 +-
src/relax/transform/allocate_workspace.cc | 19 +-
src/relax/transform/alter_op_impl.cc | 119 +-
src/relax/transform/attach_global_symbol.cc | 6 +-
src/relax/transform/bind_params.cc | 41 +-
src/relax/transform/bind_symbolic_vars.cc | 1 -
src/relax/transform/bundle_model_params.cc | 6 +-
src/relax/transform/call_tir_rewrite.cc | 55 +-
src/relax/transform/canonicalize_bindings.cc | 34 +-
src/relax/transform/combine_parallel_matmul.cc | 18 +-
src/relax/transform/compute_prim_value.cc | 2 +-
src/relax/transform/convert_layout.cc | 31 +-
src/relax/transform/dataflow_inplace.cc | 72 +-
src/relax/transform/decompose_ops.cc | 50 +-
src/relax/transform/eliminate_common_subexpr.cc | 13 +-
src/relax/transform/expand_tuple_arguments.cc | 25 +-
src/relax/transform/fold_constant.cc | 54 +-
src/relax/transform/fuse_ops.cc | 61 +-
src/relax/transform/fuse_tir.cc | 111 +-
src/relax/transform/gradient.cc | 86 +-
src/relax/transform/gradient_simplifier.cc | 16 +-
src/relax/transform/infer_amp_utils.cc | 12 +-
src/relax/transform/infer_amp_utils.h | 4 +-
src/relax/transform/infer_layout_utils.cc | 12 +-
src/relax/transform/infer_layout_utils.h | 8 +-
src/relax/transform/inline_functions.cc | 2 +-
src/relax/transform/kill_after_last_use.cc | 13 +-
src/relax/transform/lambda_lift.cc | 31 +-
src/relax/transform/lazy_transform_params.cc | 24 +-
src/relax/transform/legalize_ops.cc | 55 +-
src/relax/transform/lift_transform_params.cc | 53 +-
src/relax/transform/lower_alloc_tensor.cc | 6 +-
src/relax/transform/merge_composite_functions.cc | 6 +-
src/relax/transform/normalize.cc | 12 +-
src/relax/transform/realize_vdevice.cc | 38 +-
src/relax/transform/remove_purity_checking.cc | 11 +-
src/relax/transform/remove_unused_outputs.cc | 36 +-
src/relax/transform/remove_unused_parameters.cc | 18 +-
.../transform/reorder_permute_dims_after_concat.cc | 12 +-
src/relax/transform/reorder_take_after_matmul.cc | 33 +-
src/relax/transform/rewrite_cuda_graph.cc | 58 +-
src/relax/transform/rewrite_dataflow_reshape.cc | 20 +-
src/relax/transform/run_codegen.cc | 13 +-
.../specialize_primfunc_based_on_callsite.cc | 58 +-
src/relax/transform/split_call_tir_by_pattern.cc | 4 +-
.../transform/split_layout_rewrite_preproc.cc | 16 +-
src/relax/transform/static_plan_block_memory.cc | 40 +-
src/relax/transform/to_mixed_precision.cc | 48 +-
src/relax/transform/to_non_dataflow.cc | 3 +-
src/relax/transform/topological_sort.cc | 2 +-
...e_param_struct_info.cc => update_param_type.cc} | 24 +-
src/relax/transform/update_vdevice.cc | 11 +-
src/relax/transform/utils.cc | 30 +-
src/relax/transform/utils.h | 50 +-
src/relax/utils.cc | 71 +-
src/runtime/vm/builtin.cc | 2 +-
src/script/ir_builder/ir/ir.cc | 25 +-
src/script/printer/script_printer.cc | 4 +-
src/tirx/ir/function.cc | 31 +-
src/tirx/script/builder/ir.cc | 2 +-
tests/cpp/nested_msg_test.cc | 46 +-
.../python/codegen/test_target_codegen_cuda_fp8.py | 14 +-
.../python/contrib/test_tir_triton_integration.py | 2 +-
tests/python/disco/test_loader.py | 4 +-
tests/python/disco/test_nvshmem.py | 2 +-
tests/python/disco/test_session.py | 6 +-
tests/python/nightly/test_nnapi/infrastructure.py | 10 +-
tests/python/relax/backend/adreno/mod_utils.py | 4 +-
.../adreno/test_transform_annotate_custom_scope.py | 24 +-
.../test_transform_fold_vdevice_scope_change.py | 14 +-
tests/python/relax/backend/adreno/utils.py | 4 +-
...r_sinfo.py => test_distributed_dtensor_type.py} | 44 +-
.../test_distributed_transform_lower_distir.py | 28 +-
...ributed_transform_lower_global_to_local_view.py | 152 +-
...est_distributed_transform_propagate_sharding.py | 200 +-
.../test_distributed_tvmscript_parser.py | 10 +-
.../test_distributed_tvmscript_printer.py | 30 +-
tests/python/relax/test_analysis.py | 28 +-
.../relax/test_analysis_contains_impure_call.py | 4 +-
.../relax/test_analysis_estimate_memory_usage.py | 2 +-
..._analysis.py => test_analysis_type_analysis.py} | 460 ++--
tests/python/relax/test_analysis_well_formed.py | 236 +--
tests/python/relax/test_ast_printer.py | 127 +-
.../python/relax/test_backend_dispatch_sampling.py | 4 +-
.../relax/test_backend_transform_shape_lower.py | 126 +-
tests/python/relax/test_base_py_module_printer.py | 36 +-
.../relax/test_base_py_module_symbolic_shape.py | 20 +-
tests/python/relax/test_bind_symbolic_vars.py | 24 +-
tests/python/relax/test_blockbuilder_core.py | 198 +-
tests/python/relax/test_blockbuilder_emit_te.py | 4 +-
tests/python/relax/test_codegen_coreml.py | 44 +-
tests/python/relax/test_codegen_cutlass.py | 48 +-
tests/python/relax/test_codegen_tensorrt.py | 2 +-
tests/python/relax/test_contrib_vllm.py | 6 +-
tests/python/relax/test_dataflow_inplace.py | 34 +-
tests/python/relax/test_dataflow_pattern.py | 56 +-
tests/python/relax/test_dataflow_rewriter.py | 100 +-
...eliminate_pad_branch_using_buffer_assumption.py | 12 +-
tests/python/relax/test_expr.py | 84 +-
tests/python/relax/test_expr_functor.py | 12 +-
tests/python/relax/test_frontend_common.py | 10 +-
.../relax/test_frontend_from_exported_program.py | 32 +-
tests/python/relax/test_frontend_from_fx.py | 4 +-
.../python/relax/test_frontend_nn_extern_module.py | 4 +-
tests/python/relax/test_frontend_nn_modules.py | 6 +-
tests/python/relax/test_frontend_nn_op.py | 18 +-
tests/python/relax/test_frontend_onnx.py | 6 +-
tests/python/relax/test_frontend_tflite.py | 110 +-
tests/python/relax/test_inline_functions.py | 6 +-
tests/python/relax/test_kill_after_last_use.py | 4 +-
tests/python/relax/test_op_binary.py | 164 +-
tests/python/relax/test_op_ccl.py | 176 +-
tests/python/relax/test_op_create.py | 640 +++---
tests/python/relax/test_op_datatype.py | 72 +-
tests/python/relax/test_op_distributed.py | 4 +-
tests/python/relax/test_op_grad.py | 26 +-
tests/python/relax/test_op_gradient_numeric.py | 26 +-
tests/python/relax/test_op_image.py | 162 +-
tests/python/relax/test_op_index.py | 544 +++--
tests/python/relax/test_op_linear_algebra.py | 190 +-
tests/python/relax/test_op_manipulate.py | 2224 +++++++++-----------
tests/python/relax/test_op_misc.py | 18 +-
tests/python/relax/test_op_nn.py | 824 ++++----
tests/python/relax/test_op_nn_convolution.py | 606 +++---
tests/python/relax/test_op_nn_pooling.py | 774 +++----
tests/python/relax/test_op_qdq.py | 34 +-
tests/python/relax/test_op_sampling.py | 4 +-
tests/python/relax/test_op_search.py | 344 ++-
tests/python/relax/test_op_set.py | 554 +++--
tests/python/relax/test_op_sort.py | 150 +-
tests/python/relax/test_op_statistical.py | 284 ++-
tests/python/relax/test_op_ternary.py | 82 +-
tests/python/relax/test_op_unary.py | 74 +-
tests/python/relax/test_op_view.py | 114 +-
tests/python/relax/test_op_vision.py | 170 +-
.../python/relax/test_optimize_layout_transform.py | 34 +-
tests/python/relax/test_pipeline.py | 6 +-
tests/python/relax/test_pytorch_integration.py | 4 +-
tests/python/relax/test_relax_operators.py | 26 +-
.../python/relax/test_relax_to_pyfunc_converter.py | 14 +-
tests/python/relax/test_tir_call_source_kernel.py | 2 +-
tests/python/relax/test_training_append_loss.py | 2 +-
tests/python/relax/test_training_loss.py | 46 +-
tests/python/relax/test_training_setup_trainer.py | 14 +-
.../python/relax/test_training_trainer_numeric.py | 16 +-
tests/python/relax/test_transform.py | 30 +-
.../relax/test_transform_adjust_matmul_order.py | 10 +-
tests/python/relax/test_transform_alter_op_impl.py | 40 +-
...st_transform_attach_attr_layout_free_buffers.py | 24 +-
tests/python/relax/test_transform_bind_params.py | 22 +-
.../relax/test_transform_bind_symbolic_vars.py | 28 +-
.../relax/test_transform_canonicalize_bindings.py | 22 +-
tests/python/relax/test_transform_codegen_pass.py | 8 +-
.../relax/test_transform_compute_prim_value.py | 8 +-
tests/python/relax/test_transform_cse.py | 8 +-
.../relax/test_transform_dead_code_elimination.py | 6 +-
tests/python/relax/test_transform_decompose_ops.py | 2 +-
.../relax/test_transform_error_enrichment.py | 10 +-
tests/python/relax/test_transform_fold_constant.py | 4 +-
tests/python/relax/test_transform_fuse_ops.py | 68 +-
.../relax/test_transform_fuse_ops_by_pattern.py | 6 +-
tests/python/relax/test_transform_fuse_tir.py | 114 +-
.../relax/test_transform_fuse_transpose_matmul.py | 8 +-
.../relax/test_transform_gradient_checkpoint.py | 22 +-
.../relax/test_transform_gradient_numeric.py | 8 +-
.../relax/test_transform_gradient_te_register.py | 32 +-
.../test_transform_inline_private_functions.py | 2 +-
tests/python/relax/test_transform_lambda_lift.py | 14 +-
.../relax/test_transform_lazy_transform_params.py | 126 +-
tests/python/relax/test_transform_legalize_ops.py | 30 +-
.../relax/test_transform_legalize_ops_binary.py | 8 +-
.../relax/test_transform_legalize_ops_ccl.py | 22 +-
.../test_transform_legalize_ops_create_datatype.py | 2 +-
.../test_transform_legalize_ops_distributed.py | 4 +-
.../relax/test_transform_legalize_ops_grad.py | 14 +-
..._transform_legalize_ops_index_linear_algebra.py | 20 +-
.../test_transform_legalize_ops_manipulate.py | 44 +-
.../python/relax/test_transform_legalize_ops_nn.py | 64 +-
.../relax/test_transform_legalize_ops_qdq.py | 24 +-
...st_transform_legalize_ops_search_statistical.py | 16 +-
.../relax/test_transform_lift_transform_params.py | 24 +-
.../test_transform_lower_gpu_ipc_alloc_storage.py | 4 +-
.../test_transform_merge_composite_functions.py | 4 +-
.../relax/test_transform_meta_schedule_tuning.py | 4 +-
tests/python/relax/test_transform_normalize.py | 28 +-
.../relax/test_transform_normalize_global_var.py | 2 +-
...st_transform_operator_specific_normalization.py | 22 +-
.../relax/test_transform_rewrite_cuda_graph.py | 76 +-
.../test_transform_rewrite_dataflow_reshape.py | 42 +-
...nsform_specialize_primfunc_based_on_callsite.py | 48 +-
.../test_transform_split_layout_rewrite_preproc.py | 22 +-
.../test_transform_static_plan_block_memory.py | 14 +-
.../relax/test_transform_to_mixed_precision.py | 4 +-
...info.py => test_transform_update_param_type.py} | 10 +-
tests/python/relax/test_tvmscript_ir_builder.py | 62 +-
tests/python/relax/test_tvmscript_parser.py | 170 +-
tests/python/relax/test_tvmscript_printer_relax.py | 181 +-
.../relax/{test_struct_info.py => test_type.py} | 108 +-
tests/python/relax/test_utils.py | 2 +-
tests/python/relax/test_vm_build.py | 22 +-
tests/python/relax/test_vm_builtin.py | 2 +-
tests/python/relax/test_vm_builtin_lower.py | 12 +-
tests/python/relax/test_vm_codegen_only.py | 24 +-
tests/python/relax/test_vm_codegen_tir.py | 8 +-
tests/python/relax/test_vm_cuda_graph.py | 6 +-
tests/python/s_tir/dlight/test_benchmark.py | 12 +-
tests/python/tirx-base/test_tir_specialize.py | 20 +-
.../python/tvmscript/test_tvmscript_parser_tir.py | 62 +-
tests/python/tvmscript/test_tvmscript_roundtrip.py | 30 +-
489 files changed, 19488 insertions(+), 15367 deletions(-)
rename docs/{reference/api/python/tirx => tirx/api}/analysis.rst (100%)
copy docs/{reference/api/python/relax/frontend.rst => tirx/api/backend.rst}
(53%)
copy docs/{reference/api/python/driver.rst => tirx/api/index.rst} (83%)
rename docs/{reference/api/python/tirx => tirx/api}/stmt_functor.rst (100%)
rename docs/{reference/api/python/tirx => tirx/api}/tirx.rst (94%)
rename docs/{reference/api/python/tirx => tirx/api}/transform.rst (95%)
rename docs/{reference/api/python/tirx/backend.rst => tirx/arch/index.rst}
(85%)
create mode 100644 docs/tirx/arch/lowering_pipeline.rst
create mode 100644 docs/tirx/install.rst
create mode 100644 docs/tirx/layout.rst
create mode 100644 docs/tirx/native_basics.rst
create mode 100644 docs/tirx/native_basics/cuda/buffers.rst
create mode 100644 docs/tirx/native_basics/cuda/compiling.rst
create mode 100644 docs/tirx/native_basics/cuda/control_flow.rst
create mode 100644 docs/tirx/native_basics/cuda/data_types.rst
create mode 100644 docs/tirx/native_basics/cuda/first_kernel.rst
create mode 100644 docs/tirx/native_basics/cuda/functions.rst
create mode 100644 docs/tirx/native_basics/cuda/parser_utils.rst
create mode 100644 docs/tirx/native_basics/cuda/threads_sync.rst
create mode 100644 docs/tirx/overview.rst
create mode 100644 docs/tirx/tile_primitives.rst
create mode 100644 docs/tirx/tile_primitives/copy.rst
create mode 100644 docs/tirx/tile_primitives/copy/fallback.rst
create mode 100644 docs/tirx/tile_primitives/copy/gmem_smem.rst
create mode 100644 docs/tirx/tile_primitives/copy/ldstmatrix.rst
create mode 100644 docs/tirx/tile_primitives/copy/reg.rst
create mode 100644 docs/tirx/tile_primitives/copy_async.rst
create mode 100644 docs/tirx/tile_primitives/copy_async/dsmem.rst
create mode 100644 docs/tirx/tile_primitives/copy_async/ldgsts.rst
create mode 100644 docs/tirx/tile_primitives/copy_async/tcgen05_cp.rst
create mode 100644 docs/tirx/tile_primitives/copy_async/tcgen05_ldst.rst
create mode 100644 docs/tirx/tile_primitives/copy_async/tma.rst
create mode 100644 docs/tirx/tile_primitives/elementwise.rst
create mode 100644 docs/tirx/tile_primitives/elementwise/reg.rst
create mode 100644 docs/tirx/tile_primitives/elementwise/smem.rst
create mode 100644 docs/tirx/tile_primitives/gemm.rst
create mode 100644 docs/tirx/tile_primitives/gemm_async.rst
create mode 100644 docs/tirx/tile_primitives/permute_layout.rst
copy docs/{install/index.rst => tirx/tile_primitives/reduction.rst} (51%)
create mode 100644 docs/tirx/tile_primitives/reduction/local.rst
create mode 100644 docs/tirx/tile_primitives/reduction/shared.rst
create mode 100644 docs/tirx/tile_primitives/reduction/sm100_packed.rst
rename include/tvm/relax/distributed/{struct_info.h => type.h} (74%)
delete mode 100644 include/tvm/relax/struct_info.h
delete mode 100644 include/tvm/relax/struct_info_functor.h
rename include/tvm/{ir => relax}/type_functor.h (55%)
rename python/tvm/relax/distributed/{struct_info.py => type.py} (88%)
rename python/tvm/relax/{struct_info.py => type.py} (68%)
delete mode 100644 src/ir/type_functor.cc
rename src/relax/analysis/{struct_info_analysis.cc => type_analysis.cc} (66%)
rename src/relax/distributed/{struct_info.cc => type.cc} (81%)
create mode 100644 src/relax/ir/dependent_type.cc
delete mode 100644 src/relax/ir/struct_info.cc
delete mode 100644 src/relax/ir/struct_info_functor.cc
create mode 100644 src/relax/ir/type_functor.cc
rename src/relax/script/printer/{struct_info.cc => dependent_type.cc} (71%)
rename src/relax/transform/{update_param_struct_info.cc =>
update_param_type.cc} (78%)
rename tests/python/relax/distributed/{test_distributed_dtensor_sinfo.py =>
test_distributed_dtensor_type.py} (69%)
rename tests/python/relax/{test_analysis_struct_info_analysis.py =>
test_analysis_type_analysis.py} (58%)
rename tests/python/relax/{test_transform_update_param_struct_info.py =>
test_transform_update_param_type.py} (85%)
rename tests/python/relax/{test_struct_info.py => test_type.py} (65%)
