This is an automated email from the ASF dual-hosted git repository. ruihangl pushed a commit to branch unity-staging in repository https://gitbox.apache.org/repos/asf/tvm.git
commit 2dbab710e8fdad94092331e0efda1be91f217029 Merge: bc2415bcc8 748882aae7 Author: Ruihang Lai <[email protected]> AuthorDate: Thu Nov 16 10:30:20 2023 -0500 Merge branch 'main' into 'unity' .github/workflows/main.yml | 6 +- ci/README.md | 16 +- ci/jenkins/README.md | 4 +- cmake/utils/FindLLVM.cmake | 1 + conftest.py | 2 +- docs/contribute/ci.rst | 2 +- docs/contribute/code_guide.rst | 2 +- docs/contribute/pull_request.rst | 4 +- docs/contribute/release_process.rst | 16 +- docs/reference/api/links.rst | 1 - include/tvm/ir/expr.h | 4 +- include/tvm/runtime/threading_backend.h | 70 +++ include/tvm/target/tag.h | 11 + include/tvm/tir/stmt.h | 8 +- include/tvm/tir/var.h | 2 +- python/setup.py | 5 + python/tvm/_ffi/libinfo.py | 18 +- python/tvm/ir/expr.py | 41 +- python/tvm/relay/frontend/pytorch.py | 10 + python/tvm/relay/frontend/tflite.py | 5 - python/tvm/relay/op/nn/_nn.py | 6 + python/tvm/relay/op/strategy/adreno.py | 52 ++ .../tvm/script/_ffi_api.py | 16 +- python/tvm/script/ir_builder/tir/ir.py | 3 +- python/tvm/testing/plugin.py | 2 +- python/tvm/tir/expr.py | 379 +++++++++----- python/tvm/tir/op.py | 1 - python/tvm/tir/stmt.py | 298 ++++++++--- python/tvm/tir/tensor_intrin/cuda.py | 571 ++++++++++++++------- python/tvm/topi/adreno/__init__.py | 2 + .../tvm/topi/adreno/conv2d_transpose_alter_op.py | 121 +++++ python/tvm/topi/adreno/conv2d_transpose_nchw.py | 412 +++++++++++++++ python/tvm/topi/adreno/utils.py | 23 + python/tvm/topi/arm_cpu/arm_utils.py | 38 ++ python/tvm/topi/arm_cpu/conv2d_alter_op.py | 18 +- python/tvm/topi/nn/conv2d.py | 38 +- python/tvm/topi/nn/dense.py | 125 +++-- src/auto_scheduler/transform_step.cc | 2 +- src/relay/transforms/annotate_texture_storage.cc | 4 + src/runtime/hexagon/profiler/README.md | 2 +- src/target/tag.cc | 7 +- src/target/target_kind.cc | 1 + src/tir/ir/expr.cc | 4 +- src/tir/transforms/inject_permuted_layout.cc | 401 ++++++++------- src/topi/nn.cc | 2 +- tests/cpp/threading_backend_test.cc | 9 + tests/lint/pylint.sh | 12 +- .../test_arith_canonical_simplify.py | 0 .../test_arith_const_int_bound.py | 0 .../{unittest => arith}/test_arith_deduce_bound.py | 0 .../test_arith_detect_clip_bound.py | 0 .../{unittest => arith}/test_arith_detect_cse.py | 0 .../test_arith_detect_linear_equation.py | 0 .../test_arith_domain_touched.py | 0 .../{unittest => arith}/test_arith_intset.py | 0 .../test_arith_iter_affine_map.py | 0 .../{unittest => arith}/test_arith_modular_set.py | 0 .../test_arith_narrow_predicate_expression.py | 0 .../test_arith_rewrite_simplify.py | 0 .../{unittest => arith}/test_arith_simplify.py | 0 .../test_arith_solve_linear_equations.py | 0 .../test_arith_solve_linear_inequality.py | 0 .../test_auto_scheduler_compute_dag.py | 0 .../test_auto_scheduler_cost_model.py | 0 .../test_auto_scheduler_evolutionary_search.py | 0 .../test_auto_scheduler_feature.py | 0 .../test_auto_scheduler_layout_rewrite.py | 0 .../test_auto_scheduler_loop_state.py | 0 .../test_auto_scheduler_measure.py | 0 .../test_auto_scheduler_search_policy.py | 0 .../test_auto_scheduler_search_task.py | 0 .../test_auto_scheduler_sketch_generation.py | 0 .../test_auto_scheduler_task_scheduler.py | 0 .../{unittest => autotvm}/test_autotvm_database.py | 0 .../test_autotvm_dispatch_context.py | 0 .../test_autotvm_droplet_tuner.py | 0 .../{unittest => autotvm}/test_autotvm_feature.py | 0 .../test_autotvm_flop_calculator.py | 0 .../{unittest => autotvm}/test_autotvm_ga_tuner.py | 0 .../test_autotvm_graph_tuner_core.py | 0 .../test_autotvm_graph_tuner_utils.py | 0 .../test_autotvm_index_tuner.py | 0 .../{unittest => autotvm}/test_autotvm_measure.py | 0 .../{unittest => autotvm}/test_autotvm_record.py | 0 .../{unittest => autotvm}/test_autotvm_space.py | 0 .../test_autotvm_xgboost_model.py | 0 .../test_gpu_codegen_allreduce.py} | 0 .../{unittest => codegen}/test_inject_ptx_ldg32.py | 0 .../test_target_codegen_aarch64.py | 0 .../test_target_codegen_arm.py | 0 .../test_target_codegen_blob.py | 0 .../test_target_codegen_bool.py | 0 .../test_target_codegen_c_host.py | 0 .../test_target_codegen_cross_llvm.py | 0 .../test_target_codegen_cuda.py | 0 .../test_target_codegen_device.py | 0 .../test_target_codegen_extern.py | 0 .../test_target_codegen_hexagon.py | 0 .../test_target_codegen_llvm.py | 0 .../test_target_codegen_metal.py | 0 .../test_target_codegen_opencl.py | 0 .../test_target_codegen_rocm.py | 0 .../test_target_codegen_static_init.py | 0 .../test_target_codegen_vm_basic.py | 0 .../test_target_codegen_vulkan.py | 0 .../test_target_codegen_x86.py | 0 .../test_target_texture_codegen_opencl.py | 0 tests/python/conftest.py | 10 +- tests/python/contrib/test_random.py | 15 +- tests/python/contrib/test_rocblas.py | 8 +- tests/python/contrib/test_rpc_proxy.py | 7 +- tests/python/contrib/test_rpc_server_device.py | 1 - tests/python/contrib/test_rpc_tracker.py | 13 +- tests/python/contrib/test_sort.py | 11 +- tests/python/contrib/test_sparse.py | 12 +- tests/python/contrib/test_tedd.py | 28 +- tests/python/frontend/pytorch/test_forward.py | 51 ++ tests/python/frontend/tflite/test_forward.py | 2 +- .../test_container_structural_equal.py | 0 .../{unittest => ir}/test_datatype_nv_fp8.py | 0 tests/python/{unittest => ir}/test_ir_attrs.py | 0 tests/python/{unittest => ir}/test_ir_container.py | 0 tests/python/{unittest => ir}/test_ir_type.py | 0 .../{unittest => ir}/test_node_reflection.py | 0 tests/python/{unittest => ir}/test_object_path.py | 0 .../test_roundtrip_runtime_module.py | 0 .../test_meta_schedule_arg_info.py | 0 .../test_meta_schedule_builder.py | 0 .../test_meta_schedule_byoc_tensorrt.py | 0 .../test_meta_schedule_cost_model.py | 0 .../test_meta_schedule_cpu_dot_product.py | 0 .../test_meta_schedule_database.py | 0 .../test_meta_schedule_feature_extractor.py | 0 ...schedule_feature_extractor_per_store_feature.py | 0 .../test_meta_schedule_measure_callback.py | 0 ...meta_schedule_mma_m16n8k8_auto_tensorization.py | 0 .../test_meta_schedule_multi_anchor.py | 0 ...eta_schedule_mutator_mutate_compute_location.py | 0 .../test_meta_schedule_mutator_mutate_parallel.py | 0 ..._meta_schedule_mutator_mutate_thread_binding.py | 0 .../test_meta_schedule_mutator_mutate_tile_size.py | 0 .../test_meta_schedule_mutator_mutate_unroll.py | 0 .../test_meta_schedule_post_order_apply.py | 0 ...ule_postproc_disallow_async_strided_mem_copy.py | 0 ...meta_schedule_postproc_disallow_dynamic_loop.py | 0 ..._schedule_postproc_rewrite_cooperative_fetch.py | 0 .../test_meta_schedule_postproc_rewrite_layout.py | 0 ...e_postproc_rewrite_parallel_vectorize_unroll.py | 0 ...ta_schedule_postproc_rewrite_reduction_block.py | 0 ...est_meta_schedule_postproc_rewrite_tensorize.py | 0 ...meta_schedule_postproc_rewrite_unbound_block.py | 0 .../test_meta_schedule_postproc_verify_gpu_code.py | 0 ...est_meta_schedule_postproc_verify_vtcm_limit.py | 0 .../test_meta_schedule_profiler.py | 0 .../test_meta_schedule_relay_integration.py | 0 .../test_meta_schedule_relay_tir_compute.py | 0 .../test_meta_schedule_runner.py | 0 ...meta_schedule_schedule_cuda_layout_transform.py | 0 ...test_meta_schedule_schedule_rule_add_rfactor.py | 0 ...eta_schedule_schedule_rule_apply_custom_rule.py | 0 .../test_meta_schedule_schedule_rule_auto_bind.py | 0 ...test_meta_schedule_schedule_rule_auto_inline.py | 0 ...chedule_schedule_rule_cross_thread_reduction.py | 0 .../test_meta_schedule_schedule_rule_mlt.py | 0 .../test_meta_schedule_schedule_rule_mlt_intrin.py | 0 .../test_meta_schedule_schedule_rule_mlt_tc.py | 0 ...dule_schedule_rule_parallel_vectorize_unroll.py | 0 ...hedule_schedule_rule_random_compute_location.py | 0 .../test_meta_schedule_search_strategy.py | 0 .../test_meta_schedule_space_cpu.py | 0 .../test_meta_schedule_space_cpu_winograd.py | 0 .../test_meta_schedule_space_cuda.py | 0 .../test_meta_schedule_space_cuda_async.py | 0 ...ule_space_cuda_async_multiple_initialization.py | 0 .../test_meta_schedule_space_cuda_winograd.py | 0 .../test_meta_schedule_space_generator.py | 0 .../test_meta_schedule_task_scheduler.py | 0 .../test_meta_schedule_trace_apply.py | 0 .../test_meta_schedule_tune_context.py | 0 .../test_meta_schedule_tune_tir.py | 0 .../test_aot_legalize_packed_call.py | 0 tests/python/{unittest => micro}/test_crt.py | 0 .../test_micro_model_library_format.py | 0 .../{unittest => micro}/test_micro_ms_tuning.py | 0 .../{unittest => micro}/test_micro_project_api.py | 0 .../{unittest => micro}/test_micro_transport.py | 0 .../test_conv2d_transpose_nchw_texture.py | 325 ++++++++++++ .../relay/opencl_texture/utils/adreno_utils.py | 5 +- .../{unittest => relay}/test_custom_datatypes.py | 0 .../python/{unittest => relay}/test_link_params.py | 0 .../test_pass_div_to_mul.py} | 0 tests/python/{unittest => relay}/test_roofline.py | 0 .../test_set_input_zero_copy.py | 0 .../test_evaluator_with_preproc.py | 0 .../python/{unittest => runtime}/test_rpc_base.py | 0 .../test_runtime_container.py | 0 .../{unittest => runtime}/test_runtime_dlpack.py | 0 .../{unittest => runtime}/test_runtime_error.py | 0 .../test_runtime_extension.py | 0 .../{unittest => runtime}/test_runtime_graph.py | 0 .../test_runtime_graph_cuda_graph.py | 0 .../test_runtime_graph_debug.py | 0 .../test_runtime_heterogeneous.py | 0 .../{unittest => runtime}/test_runtime_measure.py | 0 .../test_runtime_module_based_interface.py | 0 .../test_runtime_module_export.py | 0 .../test_runtime_module_load.py | 0 .../test_runtime_module_property.py | 0 .../test_runtime_profiling.py | 0 .../{unittest => runtime}/test_runtime_rpc.py | 0 .../{unittest => runtime}/test_runtime_trace.py | 0 .../test_runtime_vm_profiler.py | 0 tests/python/{unittest => target}/test_device.py | 0 .../test_target_parser_mprofile.py | 0 .../{unittest => target}/test_target_target.py | 7 + tests/python/{unittest => te}/test_te_autodiff.py | 0 .../python/{unittest => te}/test_te_build_lower.py | 0 .../{unittest => te}/test_te_create_primfunc.py | 0 tests/python/{unittest => te}/test_te_group.py | 0 .../{unittest => te}/test_te_hybrid_script.py | 0 tests/python/{unittest => te}/test_te_schedule.py | 0 .../test_te_schedule_bound_inference.py | 0 .../test_te_schedule_bound_inference_tiling.py | 0 .../{unittest => te}/test_te_schedule_graph.py | 0 .../{unittest => te}/test_te_schedule_lstm.py | 0 .../{unittest => te}/test_te_schedule_ops.py | 0 ...te_schedule_postproc_rewrite_for_tensor_core.py | 0 .../test_te_schedule_tensor_core.py | 0 .../{unittest => te}/test_te_schedule_tensorize.py | 0 tests/python/{unittest => te}/test_te_tag.py | 0 tests/python/{unittest => te}/test_te_tensor.py | 0 .../{unittest => te}/test_te_tensor_overload.py | 0 .../{unittest => te}/test_te_transform_layout.py | 0 .../{unittest => te}/test_te_verify_compute.py | 0 .../{unittest => testing}/test_filter_untracked.py | 0 .../{unittest => testing}/test_format_si_prefix.py | 0 .../{unittest => testing}/test_gen_requirements.py | 0 tests/python/{unittest => testing}/test_testing.py | 0 .../test_tvm_testing_before_after.py | 0 .../test_tvm_testing_features.py | 0 .../test_type_annotation_checker.py | 0 ...test_tir_analysis_calculate_allocated_memory.py | 0 .../test_tir_analysis_calculate_workspace.py | 0 .../test_tir_analysis_detect_buffer_access_lca.py | 0 .../test_tir_analysis_device_constraint_utils.py} | 0 .../test_tir_analysis_estimate_tir_flops.py | 0 .../test_tir_analysis_expr_deep_equal.py | 0 .../test_tir_analysis_get_block_access_region.py | 0 .../test_tir_analysis_identify_memcpy.py | 0 .../test_tir_analysis_oob.py | 0 .../test_tir_analysis_stmt_finding.py | 0 .../test_tir_analysis_usedef.py | 0 .../test_tir_analysis_verify_gpu_code.py | 0 .../test_tir_analysis_verify_memory.py | 0 .../test_tir_analysis_verify_ssa.py | 0 .../test_tir_analysis_verify_well_formed.py | 0 tests/python/{tir => tir-base}/test_debug_info.py | 0 .../{unittest => tir-base}/test_lower_build.py | 0 .../{unittest => tir-base}/test_slice_tir.py | 0 .../python/{unittest => tir-base}/test_tir_base.py | 0 .../test_tir_block_dependence_info.py | 0 .../{unittest => tir-base}/test_tir_buffer.py | 0 .../{unittest => tir-base}/test_tir_constructor.py | 0 .../{unittest => tir-base}/test_tir_data_layout.py | 0 .../{unittest => tir-base}/test_tir_host_func.py | 0 .../{unittest => tir-base}/test_tir_imm_values.py | 0 .../{unittest => tir-base}/test_tir_index_map.py | 0 .../{unittest => tir-base}/test_tir_intrin.py | 0 .../{unittest => tir-base}/test_tir_ir_builder.py | 0 .../{unittest => tir-base}/test_tir_nodes.py | 0 .../{unittest => tir-base}/test_tir_op_types.py | 0 .../python/{unittest => tir-base}/test_tir_ops.py | 0 .../test_tir_ptx_cp_async.py | 0 .../test_tir_ptx_ldmatrix.py | 0 .../{unittest => tir-base}/test_tir_ptx_mma.py | 0 .../{unittest => tir-base}/test_tir_ptx_mma_sp.py | 0 .../{unittest => tir-base}/test_tir_renew_defs.py | 0 .../{unittest => tir-base}/test_tir_specialize.py | 0 .../test_tir_stmt_functor_ir_transform.py | 0 .../test_tir_stmt_functor_substitute.py | 0 .../test_tir_structural_equal_hash.py | 0 .../test_tir_te_extern_primfunc.py | 0 .../test_tir_texture_scope.py | 0 .../test_tir_unsafe_hide_buffer_access.py | 0 .../test_tir_schedule_analysis.py | 2 +- .../test_tir_schedule_block_scope.py | 0 .../test_tir_schedule_blockize.py | 0 .../test_tir_schedule_cache_index.py | 0 .../test_tir_schedule_cache_read_write.py | 0 .../test_tir_schedule_compute_at.py | 0 .../test_tir_schedule_compute_inline.py | 0 .../test_tir_schedule_decompose_padding.py | 0 .../test_tir_schedule_error.py | 0 .../test_tir_schedule_for_kind.py | 0 .../test_tir_schedule_instruction.py | 0 .../test_tir_schedule_merge.py | 0 .../test_tir_schedule_pad_einsum.py | 0 .../test_tir_schedule_read_write_at.py | 0 .../test_tir_schedule_reduction.py | 0 .../test_tir_schedule_reindex.py | 0 .../test_tir_schedule_reorder.py | 0 .../test_tir_schedule_reorder_block_iter_var.py} | 0 .../test_tir_schedule_rfactor.py | 0 .../test_tir_schedule_rolling_buffer.py | 0 .../test_tir_schedule_sampling.py | 0 .../test_tir_schedule_set_axis_separator.py | 0 .../test_tir_schedule_set_dtype.py | 0 .../test_tir_schedule_set_scope.py | 0 .../test_tir_schedule_split_fuse.py | 0 .../test_tir_schedule_state.py | 2 +- .../test_tir_schedule_state_cached_flags.py | 0 .../test_tir_schedule_storage_align.py | 0 .../test_tir_schedule_tensorize.py | 0 ...tir_schedule_tensorize_ldmatrix_mma_numeric.py} | 56 +- .../test_tir_schedule_tensorize_mfma_numeric.py} | 0 .../test_tir_schedule_trace.py | 0 .../test_tir_schedule_transform.py | 0 .../test_tir_schedule_transform_layout.py | 0 .../test_tir_schedule_utilities.py | 0 .../test_tir_transform_annotate_device_regions.py | 0 .../test_tir_transform_bf16_legalize.py | 0 .../test_tir_transform_combine_context_call.py | 0 .../test_tir_transform_common_subexpr_elim.py | 0 .../test_tir_transform_compact_buffer_region.py | 0 .../test_tir_transform_convert_blocks_to_opaque.py | 0 .../test_tir_transform_convert_for_loops_serial.py | 0 .../test_tir_transform_convert_ssa.py | 0 .../test_tir_transform_coproc_sync.py | 0 .../test_tir_transform_decorate_device_scope.py | 0 .../test_tir_transform_device_kernel_launch.py | 0 .../test_tir_transform_extract_constants.py | 0 .../test_tir_transform_flatten_buffer.py | 0 ...test_tir_transform_force_narrow_index_to_i32.py | 0 .../test_tir_transform_fp8_legalize.py | 0 .../test_tir_transform_helpers.py | 0 .../test_tir_transform_hoist_expression.py | 0 .../test_tir_transform_hoist_if.py | 0 .../test_tir_transform_inject_copy_intrin.py | 0 .../test_tir_transform_inject_double_buffer.py | 0 .../test_tir_transform_inject_permuted_layout.py | 351 +++++++++++++ .../test_tir_transform_inject_ptx_async_copy.py | 0 .../test_tir_transform_inject_rolling_buffer.py | 0 .../test_tir_transform_inject_software_pipeline.py | 8 +- .../test_tir_transform_inject_virtual_thread.py | 0 ...test_tir_transform_instrument_bound_checkers.py | 0 .../test_tir_transform_lift_attr_scope.py | 0 .../test_tir_transform_lift_thread_binding.py | 0 .../test_tir_transform_loop_partition.py | 0 ...t_tir_transform_lower_cross_thread_reduction.py | 0 ...r_transform_lower_device_storage_access_info.py | 0 .../test_tir_transform_lower_init_block.py | 0 .../test_tir_transform_lower_intrin.py | 0 .../test_tir_transform_lower_match_buffer.py} | 0 .../test_tir_transform_lower_opaque_block.py | 0 .../test_tir_transform_lower_thread_all_reduce.py | 0 .../test_tir_transform_lower_tvm_builtin.py | 0 .../test_tir_transform_lower_warp_memory.py | 0 .../test_tir_transform_make_packed_api.py | 0 .../test_tir_transform_make_unpacked_api.py | 0 ...transform_manifest_shared_memory_local_stage.py | 0 ...test_tir_transform_memhammer_lower_auto_copy.py | 0 ...form_merge_dynamic_shared_memory_allocations.py | 0 .../test_tir_transform_narrow_datatype.py | 0 ...sform_plan_update_buffer_allocation_location.py | 0 ...est_tir_transform_pointer_value_type_rewrite.py | 0 .../test_tir_transform_prim_func_pass.py | 0 .../test_tir_transform_profiling_instr.py | 0 ...ansform_reduce_branching_through_overcompute.py | 0 .../test_tir_transform_remove_assume.py | 0 .../test_tir_transform_remove_no_op.py | 0 .../test_tir_transform_remove_undef.py | 0 ...transform_remove_weight_layout_rewrite_block.py | 0 ...test_tir_transform_renormalize_split_pattern.py | 0 .../test_tir_transform_rewrite_unsafe_select.py | 0 .../test_tir_transform_simplify.py | 0 .../test_tir_transform_split_host_device.py | 0 .../test_tir_transform_storage_flatten.py | 0 .../test_tir_transform_storage_rewrite.py | 0 .../test_tir_transform_thread_sync.py | 0 .../test_tir_transform_unify_thread_binding.py | 0 .../test_tir_transform_unroll_loop.py | 0 .../test_tir_transform_vectorize.py | 0 .../test_transform_default_gpu_schedule.py | 0 .../{unittest => tir-usmp}/test_tir_usmp_algo.py | 0 .../test_tir_usmp_algo_hill_climb.py | 0 .../test_tir_usmp_analysis_extract_bufferinfo.py | 0 ...ransform_convert_pool_allocations_to_offsets.py | 0 .../test_tir_usmp_transform_create_io_allocates.py | 0 .../{unittest => tir-usmp}/test_tir_usmp_utils.py | 0 tests/python/topi/{python => }/common.py | 0 tests/python/topi/{python => }/test_fifo_buffer.py | 0 .../python/topi/{python => }/test_topi_argwhere.py | 0 tests/python/topi/{python => }/test_topi_basic.py | 0 .../topi/{python => }/test_topi_batch_matmul.py | 0 .../test_topi_batch_matmul_tensorcore.py | 0 .../topi/{python => }/test_topi_batch_norm.py | 0 .../{python => }/test_topi_batch_to_space_nd.py | 0 .../{python => }/test_topi_bitserial_conv2d.py | 0 .../test_topi_bitserial_conv2d_rasp.py | 0 .../topi/{python => }/test_topi_bitserial_dense.py | 0 tests/python/topi/{python => }/test_topi_bnn.py | 0 .../topi/{python => }/test_topi_broadcast.py | 0 tests/python/topi/{python => }/test_topi_clip.py | 0 tests/python/topi/{python => }/test_topi_conv1d.py | 0 .../{python => }/test_topi_conv1d_transpose_ncw.py | 0 .../topi/{python => }/test_topi_conv2d_NCHWc.py | 0 .../topi/{python => }/test_topi_conv2d_hwcn.py | 0 .../test_topi_conv2d_hwnc_tensorcore.py | 0 .../topi/{python => }/test_topi_conv2d_int8.py | 0 .../topi/{python => }/test_topi_conv2d_nchw.py | 0 .../topi/{python => }/test_topi_conv2d_nhwc.py | 0 .../test_topi_conv2d_nhwc_pack_int8.py | 0 .../test_topi_conv2d_nhwc_tensorcore.py | 0 .../{python => }/test_topi_conv2d_nhwc_winograd.py | 0 .../test_topi_conv2d_tensordot_opts.py | 0 .../test_topi_conv2d_transpose_nchw.py | 0 .../topi/{python => }/test_topi_conv2d_winograd.py | 0 .../topi/{python => }/test_topi_conv3d_ncdhw.py | 0 .../topi/{python => }/test_topi_conv3d_ndhwc.py | 0 .../test_topi_conv3d_ndhwc_tensorcore.py | 0 .../test_topi_conv3d_transpose_ncdhw.py | 0 .../topi/{python => }/test_topi_conv3d_winograd.py | 0 .../topi/{python => }/test_topi_correlation.py | 0 .../{python => }/test_topi_deformable_conv2d.py | 0 tests/python/topi/{python => }/test_topi_dense.py | 0 .../{python => }/test_topi_dense_tensorcore.py | 0 .../topi/{python => }/test_topi_depth_to_space.py | 0 .../{python => }/test_topi_depthwise_conv2d.py | 0 .../test_topi_depthwise_conv2d_back_input.py | 0 .../test_topi_depthwise_conv2d_back_weight.py | 0 tests/python/topi/{python => }/test_topi_dft.py | 0 tests/python/topi/{python => }/test_topi_dilate.py | 0 tests/python/topi/{python => }/test_topi_einsum.py | 0 .../topi/{python => }/test_topi_group_conv2d.py | 0 .../test_topi_group_conv2d_NCHWc_int8.py | 0 .../test_topi_group_conv2d_transpose.py | 0 .../topi/{python => }/test_topi_group_norm.py | 0 tests/python/topi/{python => }/test_topi_image.py | 0 .../topi/{python => }/test_topi_instance_norm.py | 0 .../topi/{python => }/test_topi_layer_norm.py | 0 tests/python/topi/{python => }/test_topi_loss.py | 0 tests/python/topi/{python => }/test_topi_lrn.py | 0 tests/python/topi/{python => }/test_topi_lstm.py | 0 tests/python/topi/{python => }/test_topi_math.py | 0 tests/python/topi/{python => }/test_topi_matmul.py | 61 ++- .../python/topi/{python => }/test_topi_pooling.py | 0 tests/python/topi/{python => }/test_topi_prng.py | 0 tests/python/topi/{python => }/test_topi_qnn.py | 0 tests/python/topi/{python => }/test_topi_reduce.py | 0 tests/python/topi/{python => }/test_topi_relu.py | 0 tests/python/topi/{python => }/test_topi_reorg.py | 0 .../python/topi/{python => }/test_topi_rms_norm.py | 0 tests/python/topi/{python => }/test_topi_scan.py | 0 .../python/topi/{python => }/test_topi_scatter.py | 0 .../topi/{python => }/test_topi_searchsorted.py | 0 .../python/topi/{python => }/test_topi_softmax.py | 0 tests/python/topi/{python => }/test_topi_sort.py | 0 .../{python => }/test_topi_space_to_batch_nd.py | 0 .../topi/{python => }/test_topi_space_to_depth.py | 0 tests/python/topi/{python => }/test_topi_sparse.py | 0 tests/python/topi/{python => }/test_topi_tensor.py | 0 .../topi/{python => }/test_topi_transform.py | 0 tests/python/topi/{python => }/test_topi_unique.py | 0 .../topi/{python => }/test_topi_upsampling.py | 0 tests/python/topi/{python => }/test_topi_util.py | 0 tests/python/topi/{python => }/test_topi_vision.py | 0 .../test_tvmscript_complete.py | 0 .../test_tvmscript_error_report.py | 0 .../test_tvmscript_ir_builder_base.py | 0 .../test_tvmscript_ir_builder_irmodule.py | 0 .../test_tvmscript_ir_builder_tir.py | 12 +- .../test_tvmscript_meta_programming.py | 0 .../{unittest => tvmscript}/test_tvmscript_ops.py | 0 .../test_tvmscript_parser_evaluator.py | 0 .../test_tvmscript_parser_ir.py | 0 .../test_tvmscript_parser_source.py | 0 .../test_tvmscript_parser_tir.py | 0 .../test_tvmscript_printer_annotation.py | 0 .../test_tvmscript_printer_doc.py | 0 .../test_tvmscript_printer_highlight.py | 0 .../test_tvmscript_printer_ir.py | 0 .../test_tvmscript_printer_metadata.py | 0 .../test_tvmscript_printer_python_doc_printer.py | 0 .../test_tvmscript_printer_structural_equal.py | 0 .../test_tvmscript_printer_tir.py | 0 .../test_tvmscript_printer_underlining.py | 0 .../test_tvmscript_regression.py | 0 .../test_tvmscript_roundtrip.py | 26 + .../test_tvmscript_syntax_sugar.py | 0 .../{unittest => tvmscript}/test_tvmscript_type.py | 0 tests/scripts/task_mypy.sh | 2 +- tests/scripts/task_python_docs.sh | 4 +- tests/scripts/task_python_unittest.sh | 28 +- tests/scripts/task_python_unittest_gpuonly.sh | 6 +- vta/runtime/runtime.cc | 12 +- 495 files changed, 2996 insertions(+), 846 deletions(-) diff --cc python/tvm/ir/expr.py index 9ca802d80e,36c425cb85..c70ac2acc7 --- a/python/tvm/ir/expr.py +++ b/python/tvm/ir/expr.py @@@ -198,47 -196,8 +215,47 @@@ class Range(Node, Scriptable) """ return _ffi_api.Range_from_min_extent(min_value, extent, span) - def __eq__(self, other): + def __eq__(self, other: Object) -> bool: return tvm.ir.structural_equal(self, other) - def __ne__(self, other): + def __ne__(self, other: Object) -> bool: return not self.__eq__(other) + + +# TODO(@relax-team): remove when we have a RelaxExpr base class +def is_relax_expr(expr: RelayExpr) -> bool: + """check if a RelayExpr is a Relax expresssion. + + Parameters + ---------- + expr : RelayExpr + The expression to check. + + Returns + ------- + res : bool + If the expression is Relax expression, return True; otherwise return False. + """ + from tvm import relax # pylint: disable=import-outside-toplevel + + if isinstance( + expr, + ( + relax.Call, + relax.Constant, + relax.Tuple, + relax.TupleGetItem, + relax.If, + relax.Var, + relax.DataflowVar, + relax.ShapeExpr, + relax.SeqExpr, + relax.Function, + relax.ExternFunc, + relax.PrimValue, + relax.StringImm, + relax.DataTypeImm, + ), + ): + return True + return False diff --cc tests/python/tir-transform/test_tir_transform_force_narrow_index_to_i32.py index c1b81853de,0000000000..c1b81853de mode 100644,000000..100644 --- a/tests/python/tir-transform/test_tir_transform_force_narrow_index_to_i32.py +++ b/tests/python/tir-transform/test_tir_transform_force_narrow_index_to_i32.py diff --cc tests/python/tir-transform/test_transform_default_gpu_schedule.py index 1af846c9d5,0000000000..1af846c9d5 mode 100644,000000..100644 --- a/tests/python/tir-transform/test_transform_default_gpu_schedule.py +++ b/tests/python/tir-transform/test_transform_default_gpu_schedule.py
