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 6f9609f10a [CI] Run s_tir/transform tests in the python-unittest stage
(#19737)
add 0ec36587af [CI] Updated cibw to 4.1.0 (#19754)
add 235387d054 [TIRX][Tests] Fix LLVM version gate for vectorized lround
(#19753)
add 126f1ba7ff [S-TIR][CUDA] Fix legacy predicated cp.async zero fill
(#19741)
add 383c4465c2 [Tests][AArch64] Make SVE codegen assertions robust across
LLVM versions (#19752)
add f5e8a3acb3 [Relax][PyTorch] Add logical_or and logical_xor converters
(#19756)
add bb6f8aec55 [TIRx] Post-bringup follow-ups: op-dispatch, namespaces,
launch bounds, gemm-async, backend reorg (#19757)
add d291513131 [REFACTOR][VM] Move CUDA graph VM builtin back under VM
runtime (#19758)
No new revisions were added by this update.
Summary of changes:
.github/actions/build-wheel-for-publish/action.yml | 2 +-
.github/workflows/publish_wheel.yml | 9 +-
CMakeLists.txt | 22 +-
.../package/manylinux_build_libtvm_runtime_cuda.sh | 24 +-
cmake/modules/CUDA.cmake | 2 +-
cmake/modules/Hexagon.cmake | 63 +-
cmake/modules/LLVM.cmake | 6 +-
cmake/modules/Metal.cmake | 4 +-
cmake/modules/OpenCL.cmake | 4 +-
cmake/modules/ROCM.cmake | 2 +-
cmake/modules/Vulkan.cmake | 14 +-
docs/reference/api/python/tirx/backend.rst | 6 -
include/tvm/tirx/builtin.h | 270 -
include/tvm/tirx/op.h | 2 -
include/tvm/tirx/target_builtin/cuda.h | 745 ---
include/tvm/tirx/target_builtin/trn.h | 156 -
python/tvm/__init__.py | 8 +-
python/tvm/backend/__init__.py | 209 +
.../default.py => backend/adreno/__init__.py} | 30 +-
.../adreno.py => backend/adreno/target_tags.py} | 2 +-
python/tvm/backend/cuda/__init__.py | 71 +
python/tvm/backend/cuda/lang/__init__.py | 70 +
.../tvm/{tirx => backend/cuda}/lang/alloc_pool.py | 10 +-
python/tvm/{tirx => backend/cuda}/lang/pipeline.py | 0
.../tvm/{tirx => backend/cuda}/lang/smem_desc.py | 2 +-
.../{tirx => backend/cuda}/lang/tile_scheduler.py | 0
.../tvm/{tirx => backend/cuda}/lang/warp_role.py | 0
python/tvm/{tirx => backend/cuda}/op.py | 4071 +--------------
.../hexagon => backend/cuda/operator}/__init__.py | 4 +-
.../cuda/operator/intrinsics}/__init__.py | 0
.../cuda}/operator/intrinsics/_schema.py | 4 +-
.../cuda/operator/intrinsics}/cp_async.py | 21 +-
.../cuda/operator/intrinsics}/header.py | 0
.../cuda/operator/intrinsics}/math.py | 4 +-
.../cuda/operator/intrinsics}/memory.py | 4 +-
.../cuda/operator/intrinsics}/misc.py | 4 +-
.../cuda/operator/intrinsics}/mma.py | 2 +-
.../cuda/operator/intrinsics}/nvshmem.py | 2 +-
.../cuda/operator/intrinsics}/registry.py | 0
.../cuda/operator/intrinsics}/sync.py | 10 +-
.../cuda/operator/intrinsics}/tcgen05.py | 2 +-
.../cuda/operator/intrinsics}/types.py | 0
.../cuda/operator/intrinsics}/utils.py | 0
.../cuda/operator/intrinsics}/wgmma.py | 2 +-
.../cuda/operator/tile_primitive}/__init__.py | 4 +
.../cuda/operator/tile_primitive}/common.py | 0
.../cuda/operator/tile_primitive}/copy/__init__.py | 0
.../cuda/operator/tile_primitive}/copy/_common.py | 0
.../operator/tile_primitive}/copy/_swizzle_iter.py | 0
.../cuda/operator/tile_primitive}/copy/fallback.py | 0
.../operator/tile_primitive}/copy/gmem_smem.py | 2 +-
.../operator/tile_primitive}/copy/ld_stmatrix.py | 2 +-
.../cuda/operator/tile_primitive}/copy/reg.py | 2 +-
.../cuda/operator/tile_primitive}/copy/utils.py | 0
.../tile_primitive}/copy_async/__init__.py | 0
.../operator/tile_primitive}/copy_async/dsmem.py | 0
.../operator/tile_primitive}/copy_async/ldgsts.py | 2 +-
.../tile_primitive}/copy_async/tcgen05_cp.py | 2 +-
.../tile_primitive}/copy_async/tcgen05_ldst.py | 0
.../operator/tile_primitive}/copy_async/tma.py | 0
.../operator/tile_primitive}/copy_async/utils.py | 0
.../tile_primitive}/elementwise/__init__.py | 2 +-
.../tile_primitive}/elementwise/_common.py | 0
.../tile_primitive}/elementwise/ops/__init__.py | 0
.../tile_primitive}/elementwise/ops/binary.py | 0
.../tile_primitive}/elementwise/ops/cast.py | 0
.../tile_primitive}/elementwise/ops/fma.py | 0
.../tile_primitive}/elementwise/ops/unary.py | 0
.../operator/tile_primitive}/elementwise/reg.py | 2 +-
.../tile_primitive}/elementwise/register.py | 0
.../operator/tile_primitive}/elementwise/smem.py | 2 +-
.../elementwise/vec_emit/__init__.py | 0
.../elementwise/vec_emit/binary_f32x2.py | 0
.../elementwise/vec_emit/cast_vec2.py | 0
.../elementwise/vec_emit/fma_f32x2.py | 0
.../operator/tile_primitive}/exec_scope_utils.py | 0
.../cuda/operator/tile_primitive}/gemm/__init__.py | 0
.../operator/tile_primitive}/gemm/mma_m16n8k_.py | 0
.../tile_primitive}/gemm_async/__init__.py | 0
.../operator/tile_primitive}/gemm_async/tcgen05.py | 211 +-
.../cuda/operator/tile_primitive}/gemm_utils.py | 0
.../cuda/operator/tile_primitive}/layout_utils.py | 0
.../tile_primitive}/permute_layout/__init__.py | 0
.../permute_layout/warp_xor_swizzle.py | 0
.../operator/tile_primitive}/reduction/__init__.py | 0
.../operator/tile_primitive}/reduction/local.py | 4 +-
.../operator/tile_primitive}/reduction/shared.py | 2 +-
.../tile_primitive}/reduction/sm100_packed.py | 2 +-
.../operator/tile_primitive}/reduction/utils.py | 9 +-
.../cuda/operator/tile_primitive}/tma_utils.py | 0
python/tvm/backend/cuda/script.py | 571 +++
.../cuda.py => backend/cuda/target_tags.py} | 2 +-
.../default.py => backend/hexagon/__init__.py} | 26 +-
.../hexagon.py => backend/hexagon/target_tags.py} | 2 +-
python/tvm/backend/metal/__init__.py | 58 +
python/tvm/backend/metal/op.py | 84 +
python/tvm/backend/metal/script.py | 55 +
.../metal.py => backend/metal/target_tags.py} | 4 +-
.../_ffi_api.py => backend/opencl/__init__.py} | 10 +-
.../_ffi_api.py => backend/rocm/__init__.py} | 10 +-
python/tvm/backend/trn/__init__.py | 68 +
python/tvm/backend/trn/layout.py | 123 +
python/tvm/backend/trn/op.py | 153 +
.../onnx => backend/trn/operator}/__init__.py | 8 +-
.../trn/operator/tile_primitive}/__init__.py | 0
.../operator/tile_primitive}/binary/__init__.py | 0
.../trn/operator/tile_primitive}/binary/default.py | 4 +-
.../trn/operator/tile_primitive}/binary/utils.py | 7 +-
.../trn/operator/tile_primitive}/common.py | 0
.../tile_primitive}/compose_op/__init__.py | 0
.../tile_primitive}/compose_op/binary_chain.py | 0
.../tile_primitive}/compose_op/binary_reduce.py | 0
.../tile_primitive}/compose_op/compose_op.py | 0
.../tile_primitive}/compose_op/reduce_negate.py | 0
.../tile_primitive}/compose_op/unary_reduce.py | 0
.../operator/tile_primitive}/compose_op/utils.py | 3 +-
.../trn/operator/tile_primitive}/copy/__init__.py | 0
.../trn/operator/tile_primitive}/copy/default.py | 11 +-
.../trn/operator/tile_primitive}/dim_utils.py | 0
.../trn/operator/tile_primitive}/gemm/__init__.py | 0
.../trn/operator/tile_primitive}/gemm/default.py | 11 +-
.../tile_primitive}/instruction_generator.py | 11 +-
.../trn/operator/tile_primitive}/private_alloc.py | 6 +-
.../operator/tile_primitive}/reduction/__init__.py | 0
.../operator/tile_primitive}/reduction/default.py | 2 +-
.../operator/tile_primitive}/reduction/utils.py | 11 +-
.../operator/tile_primitive}/select/__init__.py | 0
.../trn/operator/tile_primitive}/select/default.py | 5 +-
.../trn/operator/tile_primitive}/unary/__init__.py | 0
.../trn/operator/tile_primitive}/unary/default.py | 4 +-
.../trn/operator/tile_primitive}/unary/utils.py | 9 +-
.../tile_primitive}/unary/with_bias_scale.py | 4 +-
.../operator/tile_primitive}/workspace_utils.py | 0
python/tvm/backend/trn/pipeline.py | 58 +
python/tvm/backend/trn/script.py | 58 +
.../trn/target_tags.py} | 26 +-
.../trn => backend/trn/transform}/__init__.py | 27 +-
.../trn/transform}/naive_allocator.py | 3 +-
.../trn/transform}/private_buffer_alloc.py | 0
.../_ffi_api.py => backend/vulkan/__init__.py} | 10 +-
.../_ffi_api.py => backend/webgpu/__init__.py} | 10 +-
.../frontend/torch/base_fx_graph_translator.py | 22 +
.../frontend/torch/exported_program_translator.py | 2 +
python/tvm/relax/frontend/torch/fx_translator.py | 2 +
python/tvm/s_tir/tensor_intrin/cuda.py | 4 +-
python/tvm/s_tir/tensor_intrin/metal.py | 8 +-
python/tvm/target/tag_registry/__init__.py | 4 -
python/tvm/tirx/__init__.py | 17 +-
python/tvm/tirx/backend/__init__.py | 4 +-
python/tvm/tirx/backend/adreno/__init__.py | 17 -
python/tvm/tirx/bench.py | 8 +-
python/tvm/tirx/compilation_pipeline.py | 39 +-
python/tvm/tirx/lang/alloc_pool.py | 527 +-
python/tvm/tirx/lang/pipeline.py | 242 +-
python/tvm/tirx/lang/smem_desc.py | 53 +-
python/tvm/tirx/lang/tile_scheduler.py | 814 +--
python/tvm/tirx/lang/warp_role.py | 142 +-
python/tvm/tirx/op.py | 5285 +-------------------
python/tvm/tirx/operator/intrinsics/_common.py | 2 +-
.../tvm/tirx/operator/tile_primitive/__init__.py | 11 +-
.../operator/tile_primitive/dispatch_context.py | 4 +
python/tvm/tirx/script/builder/ir.py | 656 +--
python/tvm/tirx/transform/__init__.py | 1 -
python/tvm/tirx/transform/transform.py | 11 +
src/arith/const_int_bound.cc | 7 +-
src/arith/ir_mutator_with_analyzer.cc | 9 +-
src/arith/ir_visitor_with_analyzer.cc | 5 +-
src/arith/rewrite_simplify.cc | 11 +-
.../cuda => backend/cuda/codegen}/codegen_cuda.cc | 74 +-
.../cuda => backend/cuda/codegen}/codegen_cuda.h | 2 +-
.../cuda/codegen}/cuda_fallback_module.cc | 2 +-
.../cuda/codegen}/cuda_fallback_module.h | 4 +-
.../cuda/codegen}/intrin_rule_cuda.cc | 49 +-
.../cuda/codegen}/literal/cuda_half_t.h | 0
.../cuda/codegen}/literal/cuda_int8_t.h | 0
.../cuda/codegen}/llvm/codegen_nvptx.cc | 14 +-
src/{target/cuda => backend/cuda/codegen}/ptx.cc | 28 +-
src/{target/cuda => backend/cuda/codegen}/ptx.h | 0
src/backend/cuda/codegen/register.cc | 143 +
.../backend/cuda/op/register.cc | 24 +-
.../cuda.cc => backend/cuda/op/target_builtin.cc} | 69 +-
.../cuda => backend/cuda/runtime}/cuda_common.h | 2 +-
.../cuda/runtime}/cuda_device_api.cc | 0
.../cuda => backend/cuda/runtime}/cuda_module.cc | 8 +-
.../cuda/runtime}/l2_cache_flush.cc | 2 +-
.../hexagon/codegen}/hexagon_fallback_module.cc | 4 +-
.../hexagon/codegen}/hexagon_fallback_module.h | 4 +-
.../hexagon/codegen}/llvm/codegen_hexagon.cc | 16 +-
.../hexagon/codegen}/llvm/intrin_rule_hexagon.cc | 35 +-
src/backend/hexagon/codegen/register.cc | 64 +
.../hexagon => backend/hexagon/runtime}/README.md | 0
.../hexagon/runtime}/hexagon_buffer.cc | 0
.../hexagon/runtime}/hexagon_buffer.h | 0
.../hexagon/runtime}/hexagon_buffer_manager.h | 0
.../hexagon/runtime}/hexagon_common.cc | 0
.../hexagon/runtime}/hexagon_common.h | 0
.../hexagon/runtime}/hexagon_device_api.cc | 2 +-
.../hexagon/runtime}/hexagon_device_api.h | 0
.../hexagon/runtime}/hexagon_htp.cc | 0
.../hexagon/runtime}/hexagon_htp.h | 0
.../hexagon/runtime}/hexagon_hvx.cc | 0
.../hexagon/runtime}/hexagon_hvx.h | 0
.../hexagon/runtime}/hexagon_module.cc | 10 +-
.../hexagon/runtime}/hexagon_power_manager.cc | 0
.../hexagon/runtime}/hexagon_power_manager.h | 0
.../hexagon/runtime}/hexagon_thread_manager.cc | 0
.../hexagon/runtime}/hexagon_thread_manager.h | 0
.../hexagon/runtime}/hexagon_user_dma.cc | 0
.../hexagon/runtime}/hexagon_user_dma.h | 0
.../runtime}/hexagon_user_dma_descriptors.h | 0
.../runtime}/hexagon_user_dma_instructions.h | 0
.../hexagon/runtime}/hexagon_user_dma_registers.h | 0
.../hexagon/runtime}/hexagon_vtcm_pool.cc | 0
.../hexagon/runtime}/hexagon_vtcm_pool.h | 0
.../hexagon/runtime}/ops/conv2d.h | 0
.../hexagon/runtime}/ops/conv2d_fp16_hvx.cc | 0
.../hexagon/runtime}/ops/conv2d_quant_hvx.cc | 0
.../hexagon/runtime}/ops/conv_utils.cc | 0
.../hexagon/runtime}/profiler/README.md | 0
.../hexagon/runtime}/profiler/lwp_handler.S | 0
.../hexagon/runtime}/profiler/prof_utils.cc | 0
.../hexagon/runtime}/profiler/prof_utils.h | 0
.../hexagon/runtime}/qhl/qhl_wrapper.cc | 0
.../hexagon/runtime}/ring_buffer.h | 0
.../hexagon/runtime}/rpc/android/session.cc | 6 +-
.../hexagon/runtime}/rpc/android_bash.sh.template | 0
.../hexagon/runtime}/rpc/hexagon/rpc_server.cc | 6 +-
.../hexagon/runtime}/rpc/hexagon_rpc.idl | 0
.../runtime}/rpc/simulator/hexagon_sim_proto.h | 0
.../hexagon/runtime}/rpc/simulator/rpc_server.cc | 2 +-
.../hexagon/runtime}/rpc/simulator/session.cc | 6 +-
.../metal/codegen}/codegen_metal.cc | 24 +-
.../metal/codegen}/codegen_metal.h | 2 +-
.../metal/codegen}/intrin_rule_metal.cc | 22 +-
.../metal/codegen}/metal_fallback_module.cc | 2 +-
.../metal/codegen}/metal_fallback_module.h | 4 +-
src/backend/metal/codegen/register.cc | 63 +
.../backend/metal/op/register.cc | 24 +-
src/backend/metal/op/target_builtin.cc | 62 +
.../metal => backend/metal/runtime}/metal_common.h | 2 +-
.../metal/runtime}/metal_device_api.mm | 0
.../metal/runtime}/metal_module.mm | 10 +-
.../opencl/codegen}/codegen_opencl.cc | 18 +-
.../opencl/codegen}/codegen_opencl.h | 2 +-
.../opencl/codegen}/intrin_rule_opencl.cc | 38 +-
.../opencl/codegen}/opencl_fallback_module.cc | 4 +-
.../opencl/codegen}/opencl_fallback_module.h | 4 +-
src/backend/opencl/codegen/register.cc | 67 +
.../opencl/runtime}/opencl_common.h | 8 +-
.../opencl/runtime}/opencl_device_api.cc | 4 +-
.../opencl/runtime}/opencl_module.cc | 13 +-
.../opencl/runtime}/opencl_wrapper/README.md | 0
.../runtime}/opencl_wrapper/opencl_wrapper.cc | 0
.../opencl/runtime}/source_utils.h | 0
.../opencl => backend/opencl/runtime}/texture.h | 0
.../rocm/codegen}/llvm/codegen_amdgpu.cc | 14 +-
.../rocm/codegen}/llvm/intrin_rule_rocm.cc | 8 +-
src/backend/rocm/codegen/register.cc | 148 +
.../rocm/codegen}/rocm_fallback_module.cc | 2 +-
.../rocm/codegen}/rocm_fallback_module.h | 4 +-
.../rocm => backend/rocm/runtime}/rocm_common.h | 2 +-
.../rocm/runtime}/rocm_device_api.cc | 0
.../rocm => backend/rocm/runtime}/rocm_module.cc | 8 +-
.../source => backend/trn/codegen}/codegen_trn.cc | 59 +-
.../source => backend/trn/codegen}/codegen_trn.h | 2 +-
.../trn/codegen/register.cc} | 59 +-
.../backend/trn/op/register.cc | 24 +-
.../trn.cc => backend/trn/op/target_builtin.cc} | 34 +-
.../trn/transform/lower_trainium_layout.cc} | 91 +-
.../vulkan/codegen}/build_vulkan.cc | 14 +-
.../vulkan/codegen}/codegen_spirv.cc | 21 +-
.../vulkan/codegen}/codegen_spirv.h | 4 +-
.../vulkan/codegen}/intrin_rule_spirv.cc | 26 +-
.../vulkan/codegen}/ir_builder.cc | 0
.../vulkan => backend/vulkan/codegen}/ir_builder.h | 0
src/backend/vulkan/codegen/register.cc | 93 +
.../vulkan/codegen}/spirv_support.cc | 0
.../vulkan/codegen}/spirv_support.h | 0
.../vulkan/codegen}/spirv_utils.cc | 4 +-
.../vulkan/codegen}/spirv_utils.h | 2 +-
.../vulkan/codegen}/vulkan_fallback_module.cc | 8 +-
.../vulkan/codegen}/vulkan_fallback_module.h | 4 +-
.../vulkan => backend/vulkan/runtime}/README.md | 0
.../vulkan/runtime}/spirv_shader.h | 0
.../vulkan => backend/vulkan/runtime}/thread_map.h | 0
.../vulkan/runtime}/vulkan_amdrgp.cc | 0
.../vulkan/runtime}/vulkan_amdrgp.h | 0
.../vulkan/runtime}/vulkan_buffer.cc | 0
.../vulkan/runtime}/vulkan_buffer.h | 0
.../vulkan/runtime}/vulkan_common.cc | 0
.../vulkan/runtime}/vulkan_common.h | 0
.../vulkan/runtime}/vulkan_device.cc | 2 +-
.../vulkan/runtime}/vulkan_device.h | 0
.../vulkan/runtime}/vulkan_device_api.cc | 0
.../vulkan/runtime}/vulkan_device_api.h | 2 +-
.../vulkan/runtime}/vulkan_instance.cc | 2 +-
.../vulkan/runtime}/vulkan_instance.h | 0
.../vulkan/runtime}/vulkan_module.cc | 8 +-
.../vulkan/runtime}/vulkan_stream.cc | 2 +-
.../vulkan/runtime}/vulkan_stream.h | 0
.../vulkan/runtime}/vulkan_wrapped_func.cc | 4 +-
.../vulkan/runtime}/vulkan_wrapped_func.h | 6 +-
.../webgpu/codegen}/codegen_webgpu.cc | 20 +-
.../webgpu/codegen}/codegen_webgpu.h | 2 +-
.../webgpu/codegen}/intrin_rule_webgpu.cc | 20 +-
src/backend/webgpu/codegen/register.cc | 80 +
.../webgpu/codegen}/webgpu_fallback_module.cc | 2 +-
.../webgpu/codegen}/webgpu_fallback_module.h | 4 +-
src/relax/transform/static_plan_block_memory.cc | 2 +-
src/runtime/extra/contrib/clml/clml_runtime.h | 2 +-
.../extra/contrib/cublas/cublas_json_runtime.cc | 2 +-
src/runtime/extra/contrib/cublas/cublas_utils.cc | 2 +-
.../contrib/cudnn/cudnn_frontend/attention.cc | 2 +-
src/runtime/extra/contrib/cudnn/cudnn_utils.h | 2 +-
src/runtime/extra/contrib/curand/curand.cc | 2 +-
.../cutlass/fp16_group_gemm_runner_sm100.cuh | 2 +-
.../cutlass/fp16_group_gemm_runner_sm90.cuh | 2 +-
.../fp8_groupwise_scaled_gemm_runner_sm100.cuh | 2 +-
.../fp8_groupwise_scaled_gemm_runner_sm90.cuh | 2 +-
...p8_groupwise_scaled_group_gemm_runner_sm100.cuh | 2 +-
src/runtime/extra/contrib/cutlass/gemm_runner.cuh | 2 +-
.../extra/contrib/hipblas/hipblas_json_runtime.cc | 2 +-
src/runtime/extra/contrib/hipblas/hipblas_utils.cc | 2 +-
src/runtime/extra/contrib/nvshmem/dist_gemm.cu | 2 +-
src/runtime/extra/contrib/nvshmem/init.cc | 2 +-
.../extra/contrib/nvshmem/memory_allocator.cc | 2 +-
.../extra/contrib/tensorrt/tensorrt_calibrator.h | 2 +-
src/runtime/extra/contrib/thrust/thrust.cu | 2 +-
.../extra/disco/cuda_ipc/cuda_ipc_memory.cc | 2 +-
src/runtime/extra/disco/nccl/nccl_context.h | 4 +-
src/runtime/vm/attn_utils.h | 2 +-
src/runtime/vm/cuda/cuda_graph_builtin.cc | 4 +-
src/runtime/vm/hexagon/builtin.cc | 2 +-
src/s_tir/backend/adreno/inject_texture_alloc.cc | 2 +-
src/s_tir/backend/adreno/texture_flatten.cc | 2 +-
.../postproc/rewrite_cooperative_fetch.cc | 7 +-
src/s_tir/schedule/analysis/analysis.cc | 5 +-
src/s_tir/transform/inject_permuted_layout.cc | 29 +-
src/s_tir/transform/inject_ptx_async_copy.cc | 10 +-
src/s_tir/transform/inject_ptx_ldg32.cc | 4 +-
src/s_tir/transform/inject_software_pipeline.cc | 24 +-
src/s_tir/transform/memhammer_lower_auto_copy.cc | 6 +-
.../transform/memhammer_tensorcore_rewrite.cc | 7 +-
.../transform/merge_shared_memory_allocations.cc | 4 +-
src/s_tir/transform/tensorcore_infer_fragment.cc | 13 +-
src/target/llvm/intrin_rule_nvptx.cc | 3 +-
src/target/tag.cc | 13 -
src/target/target_kind.cc | 318 --
src/tirx/analysis/filter_canonical.cc | 5 +-
src/tirx/ir/data_type_rewriter.cc | 9 +-
src/tirx/ir/exec_scope.cc | 4 +-
src/tirx/ir/stmt.cc | 5 +-
src/tirx/op/builtin.cc | 153 -
src/tirx/op/op.cc | 47 +-
src/tirx/script/builder/frame.cc | 3 +-
src/tirx/transform/common_subexpr_elim.cc | 26 +-
src/tirx/transform/lower_intrin.cc | 4 +-
src/tirx/transform/lower_tvm_builtin.cc | 24 +-
src/tirx/transform/lower_warp_memory.cc | 51 +-
src/tirx/transform/remove_no_op.cc | 11 +-
src/tirx/transform/tile_primitive_dispatch.cc | 5 +-
.../python/codegen/test_target_codegen_aarch64.py | 22 +-
.../relax/test_frontend_from_exported_program.py | 56 +
tests/python/relax/test_frontend_from_fx.py | 50 +
.../s_tir/dlight/test_gpu_matmul_tensorize.py | 20 +-
.../test_s_tir_transform_inject_ptx_async_copy.py | 211 +-
tests/python/tirx-base/test_tir_op_types.py | 15 +-
.../tirx-transform/test_tir_transform_vectorize.py | 5 +-
.../tile_primitive/cuda/copy/test_fallback.py | 4 +-
.../tile_primitive/cuda/copy/test_gmem_smem.py | 2 +-
.../tile_primitive/cuda/copy/test_swizzle_iter.py | 10 +-
.../tile_primitive/cuda/copy_async/test_dsmem.py | 2 +-
.../cuda/copy_async/test_smem_tmem.py | 2 +-
.../tile_primitive/cuda/copy_async/test_tma.py | 8 +-
.../tile_primitive/cuda/elementwise/test_unary.py | 10 +-
.../cuda/gemm_async/test_gemm_async.py | 92 +-
.../cuda/permute_layout/test_permute_layout.py | 4 +-
.../tile_primitive/trn/test_compose_op_trn.py | 14 +-
.../operator/tile_primitive/trn/test_copy_trn.py | 8 +-
.../operator/tile_primitive/trn/test_gemm_trn.py | 4 +-
.../tile_primitive/trn/test_private_alloc_trn.py | 2 +-
.../tile_primitive/trn/test_reduction_trn.py | 4 +-
.../operator/tile_primitive/trn/test_unary_trn.py | 2 +-
tests/python/tirx/test_alloc_pool.py | 6 +-
tests/python/tirx/test_layout.py | 10 +-
tests/python/tirx/test_op_namespace_cleanup.py | 121 +
tests/python/tirx/test_printer_tir_namespaces.py | 198 +-
.../transform/test_transform_naive_allocator.py | 2 +-
388 files changed, 4488 insertions(+), 14450 deletions(-)
delete mode 100644 include/tvm/tirx/target_builtin/cuda.h
delete mode 100644 include/tvm/tirx/target_builtin/trn.h
create mode 100644 python/tvm/backend/__init__.py
copy python/tvm/{tirx/operator/tile_primitive/trn/reduction/default.py =>
backend/adreno/__init__.py} (59%)
rename python/tvm/{target/tag_registry/adreno.py =>
backend/adreno/target_tags.py} (97%)
create mode 100644 python/tvm/backend/cuda/__init__.py
create mode 100644 python/tvm/backend/cuda/lang/__init__.py
copy python/tvm/{tirx => backend/cuda}/lang/alloc_pool.py (97%)
copy python/tvm/{tirx => backend/cuda}/lang/pipeline.py (100%)
copy python/tvm/{tirx => backend/cuda}/lang/smem_desc.py (95%)
copy python/tvm/{tirx => backend/cuda}/lang/tile_scheduler.py (100%)
copy python/tvm/{tirx => backend/cuda}/lang/warp_role.py (100%)
copy python/tvm/{tirx => backend/cuda}/op.py (56%)
copy python/tvm/{contrib/hexagon => backend/cuda/operator}/__init__.py (88%)
rename python/tvm/{tirx/operator/intrinsics/cuda =>
backend/cuda/operator/intrinsics}/__init__.py (100%)
rename python/tvm/{tirx => backend/cuda}/operator/intrinsics/_schema.py (97%)
rename python/tvm/{tirx/operator/intrinsics/cuda =>
backend/cuda/operator/intrinsics}/cp_async.py (97%)
rename python/tvm/{tirx/operator/intrinsics/cuda =>
backend/cuda/operator/intrinsics}/header.py (100%)
rename python/tvm/{tirx/operator/intrinsics/cuda =>
backend/cuda/operator/intrinsics}/math.py (99%)
rename python/tvm/{tirx/operator/intrinsics/cuda =>
backend/cuda/operator/intrinsics}/memory.py (99%)
rename python/tvm/{tirx/operator/intrinsics/cuda =>
backend/cuda/operator/intrinsics}/misc.py (99%)
rename python/tvm/{tirx/operator/intrinsics/cuda =>
backend/cuda/operator/intrinsics}/mma.py (99%)
rename python/tvm/{tirx/operator/intrinsics/cuda =>
backend/cuda/operator/intrinsics}/nvshmem.py (99%)
rename python/tvm/{tirx/operator/intrinsics/cuda =>
backend/cuda/operator/intrinsics}/registry.py (100%)
rename python/tvm/{tirx/operator/intrinsics/cuda =>
backend/cuda/operator/intrinsics}/sync.py (99%)
rename python/tvm/{tirx/operator/intrinsics/cuda =>
backend/cuda/operator/intrinsics}/tcgen05.py (99%)
rename python/tvm/{tirx/operator/intrinsics/cuda =>
backend/cuda/operator/intrinsics}/types.py (100%)
rename python/tvm/{tirx/operator/intrinsics/cuda =>
backend/cuda/operator/intrinsics}/utils.py (100%)
rename python/tvm/{tirx/operator/intrinsics/cuda =>
backend/cuda/operator/intrinsics}/wgmma.py (99%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/__init__.py (89%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/common.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/copy/__init__.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/copy/_common.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/copy/_swizzle_iter.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/copy/fallback.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/copy/gmem_smem.py (99%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/copy/ld_stmatrix.py (99%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/copy/reg.py (99%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/copy/utils.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/copy_async/__init__.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/copy_async/dsmem.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/copy_async/ldgsts.py (99%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/copy_async/tcgen05_cp.py (99%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/copy_async/tcgen05_ldst.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/copy_async/tma.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/copy_async/utils.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/elementwise/__init__.py (95%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/elementwise/_common.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/elementwise/ops/__init__.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/elementwise/ops/binary.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/elementwise/ops/cast.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/elementwise/ops/fma.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/elementwise/ops/unary.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/elementwise/reg.py (99%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/elementwise/register.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/elementwise/smem.py (99%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/elementwise/vec_emit/__init__.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/elementwise/vec_emit/binary_f32x2.py
(100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/elementwise/vec_emit/cast_vec2.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/elementwise/vec_emit/fma_f32x2.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/exec_scope_utils.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/gemm/__init__.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/gemm/mma_m16n8k_.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/gemm_async/__init__.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/gemm_async/tcgen05.py (83%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/gemm_utils.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/layout_utils.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/permute_layout/__init__.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/permute_layout/warp_xor_swizzle.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/reduction/__init__.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/reduction/local.py (99%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/reduction/shared.py (99%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/reduction/sm100_packed.py (99%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/reduction/utils.py (97%)
rename python/tvm/{tirx/operator/tile_primitive/cuda =>
backend/cuda/operator/tile_primitive}/tma_utils.py (100%)
create mode 100644 python/tvm/backend/cuda/script.py
rename python/tvm/{target/tag_registry/cuda.py => backend/cuda/target_tags.py}
(99%)
copy python/tvm/{tirx/operator/tile_primitive/trn/reduction/default.py =>
backend/hexagon/__init__.py} (65%)
rename python/tvm/{target/tag_registry/hexagon.py =>
backend/hexagon/target_tags.py} (98%)
create mode 100644 python/tvm/backend/metal/__init__.py
create mode 100644 python/tvm/backend/metal/op.py
create mode 100644 python/tvm/backend/metal/script.py
rename python/tvm/{target/tag_registry/metal.py =>
backend/metal/target_tags.py} (94%)
copy python/tvm/{s_tir/backend/adreno/transform/_ffi_api.py =>
backend/opencl/__init__.py} (83%)
copy python/tvm/{s_tir/backend/adreno/transform/_ffi_api.py =>
backend/rocm/__init__.py} (83%)
create mode 100644 python/tvm/backend/trn/__init__.py
create mode 100644 python/tvm/backend/trn/layout.py
create mode 100644 python/tvm/backend/trn/op.py
copy python/tvm/{relax/frontend/onnx => backend/trn/operator}/__init__.py (87%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/__init__.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/binary/__init__.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/binary/default.py (97%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/binary/utils.py (97%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/common.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/compose_op/__init__.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/compose_op/binary_chain.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/compose_op/binary_reduce.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/compose_op/compose_op.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/compose_op/reduce_negate.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/compose_op/unary_reduce.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/compose_op/utils.py (95%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/copy/__init__.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/copy/default.py (96%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/dim_utils.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/gemm/__init__.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/gemm/default.py (97%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/instruction_generator.py (98%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/private_alloc.py (96%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/reduction/__init__.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/reduction/default.py (94%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/reduction/utils.py (95%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/select/__init__.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/select/default.py (97%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/unary/__init__.py (100%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/unary/default.py (96%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/unary/utils.py (96%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/unary/with_bias_scale.py (96%)
rename python/tvm/{tirx/operator/tile_primitive/trn =>
backend/trn/operator/tile_primitive}/workspace_utils.py (100%)
create mode 100644 python/tvm/backend/trn/pipeline.py
create mode 100644 python/tvm/backend/trn/script.py
copy python/tvm/{s_tir/meta_schedule/mutator/mutate_unroll.py =>
backend/trn/target_tags.py} (62%)
rename python/tvm/{tirx/transform/trn => backend/trn/transform}/__init__.py
(68%)
rename python/tvm/{tirx/transform/trn =>
backend/trn/transform}/naive_allocator.py (98%)
rename python/tvm/{tirx/transform/trn =>
backend/trn/transform}/private_buffer_alloc.py (100%)
copy python/tvm/{s_tir/backend/adreno/transform/_ffi_api.py =>
backend/vulkan/__init__.py} (83%)
copy python/tvm/{s_tir/backend/adreno/transform/_ffi_api.py =>
backend/webgpu/__init__.py} (83%)
delete mode 100644 python/tvm/tirx/backend/adreno/__init__.py
rename src/{target/cuda => backend/cuda/codegen}/codegen_cuda.cc (96%)
rename src/{target/cuda => backend/cuda/codegen}/codegen_cuda.h (99%)
rename src/{target/cuda => backend/cuda/codegen}/cuda_fallback_module.cc (99%)
rename src/{target/cuda => backend/cuda/codegen}/cuda_fallback_module.h (98%)
rename src/{target/cuda => backend/cuda/codegen}/intrin_rule_cuda.cc (89%)
rename src/{target/cuda => backend/cuda/codegen}/literal/cuda_half_t.h (100%)
rename src/{target/cuda => backend/cuda/codegen}/literal/cuda_int8_t.h (100%)
rename src/{target/cuda => backend/cuda/codegen}/llvm/codegen_nvptx.cc (97%)
rename src/{target/cuda => backend/cuda/codegen}/ptx.cc (97%)
rename src/{target/cuda => backend/cuda/codegen}/ptx.h (100%)
create mode 100644 src/backend/cuda/codegen/register.cc
copy tests/cpp/topi_ewise_test.cc => src/backend/cuda/op/register.cc (71%)
rename src/{tirx/op/target_builtin/cuda.cc =>
backend/cuda/op/target_builtin.cc} (94%)
rename src/{runtime/cuda => backend/cuda/runtime}/cuda_common.h (98%)
rename src/{runtime/cuda => backend/cuda/runtime}/cuda_device_api.cc (100%)
rename src/{runtime/cuda => backend/cuda/runtime}/cuda_module.cc (98%)
rename src/{runtime/cuda => backend/cuda/runtime}/l2_cache_flush.cc (96%)
rename src/{target/hexagon =>
backend/hexagon/codegen}/hexagon_fallback_module.cc (97%)
rename src/{target/hexagon =>
backend/hexagon/codegen}/hexagon_fallback_module.h (98%)
rename src/{target/hexagon => backend/hexagon/codegen}/llvm/codegen_hexagon.cc
(98%)
rename src/{target/hexagon =>
backend/hexagon/codegen}/llvm/intrin_rule_hexagon.cc (89%)
create mode 100644 src/backend/hexagon/codegen/register.cc
rename src/{runtime/hexagon => backend/hexagon/runtime}/README.md (100%)
rename src/{runtime/hexagon => backend/hexagon/runtime}/hexagon_buffer.cc
(100%)
rename src/{runtime/hexagon => backend/hexagon/runtime}/hexagon_buffer.h (100%)
rename src/{runtime/hexagon =>
backend/hexagon/runtime}/hexagon_buffer_manager.h (100%)
rename src/{runtime/hexagon => backend/hexagon/runtime}/hexagon_common.cc
(100%)
rename src/{runtime/hexagon => backend/hexagon/runtime}/hexagon_common.h (100%)
rename src/{runtime/hexagon => backend/hexagon/runtime}/hexagon_device_api.cc
(99%)
rename src/{runtime/hexagon => backend/hexagon/runtime}/hexagon_device_api.h
(100%)
rename src/{runtime/hexagon => backend/hexagon/runtime}/hexagon_htp.cc (100%)
rename src/{runtime/hexagon => backend/hexagon/runtime}/hexagon_htp.h (100%)
rename src/{runtime/hexagon => backend/hexagon/runtime}/hexagon_hvx.cc (100%)
rename src/{runtime/hexagon => backend/hexagon/runtime}/hexagon_hvx.h (100%)
rename src/{runtime/hexagon => backend/hexagon/runtime}/hexagon_module.cc (94%)
rename src/{runtime/hexagon =>
backend/hexagon/runtime}/hexagon_power_manager.cc (100%)
rename src/{runtime/hexagon =>
backend/hexagon/runtime}/hexagon_power_manager.h (100%)
rename src/{runtime/hexagon =>
backend/hexagon/runtime}/hexagon_thread_manager.cc (100%)
rename src/{runtime/hexagon =>
backend/hexagon/runtime}/hexagon_thread_manager.h (100%)
rename src/{runtime/hexagon => backend/hexagon/runtime}/hexagon_user_dma.cc
(100%)
rename src/{runtime/hexagon => backend/hexagon/runtime}/hexagon_user_dma.h
(100%)
rename src/{runtime/hexagon =>
backend/hexagon/runtime}/hexagon_user_dma_descriptors.h (100%)
rename src/{runtime/hexagon =>
backend/hexagon/runtime}/hexagon_user_dma_instructions.h (100%)
rename src/{runtime/hexagon =>
backend/hexagon/runtime}/hexagon_user_dma_registers.h (100%)
rename src/{runtime/hexagon => backend/hexagon/runtime}/hexagon_vtcm_pool.cc
(100%)
rename src/{runtime/hexagon => backend/hexagon/runtime}/hexagon_vtcm_pool.h
(100%)
rename src/{runtime/hexagon => backend/hexagon/runtime}/ops/conv2d.h (100%)
rename src/{runtime/hexagon => backend/hexagon/runtime}/ops/conv2d_fp16_hvx.cc
(100%)
rename src/{runtime/hexagon =>
backend/hexagon/runtime}/ops/conv2d_quant_hvx.cc (100%)
rename src/{runtime/hexagon => backend/hexagon/runtime}/ops/conv_utils.cc
(100%)
rename src/{runtime/hexagon => backend/hexagon/runtime}/profiler/README.md
(100%)
rename src/{runtime/hexagon => backend/hexagon/runtime}/profiler/lwp_handler.S
(100%)
rename src/{runtime/hexagon => backend/hexagon/runtime}/profiler/prof_utils.cc
(100%)
rename src/{runtime/hexagon => backend/hexagon/runtime}/profiler/prof_utils.h
(100%)
rename src/{runtime/hexagon => backend/hexagon/runtime}/qhl/qhl_wrapper.cc
(100%)
rename src/{runtime/hexagon => backend/hexagon/runtime}/ring_buffer.h (100%)
rename src/{runtime/hexagon => backend/hexagon/runtime}/rpc/android/session.cc
(96%)
rename src/{runtime/hexagon =>
backend/hexagon/runtime}/rpc/android_bash.sh.template (100%)
rename src/{runtime/hexagon =>
backend/hexagon/runtime}/rpc/hexagon/rpc_server.cc (98%)
rename src/{runtime/hexagon => backend/hexagon/runtime}/rpc/hexagon_rpc.idl
(100%)
rename src/{runtime/hexagon =>
backend/hexagon/runtime}/rpc/simulator/hexagon_sim_proto.h (100%)
rename src/{runtime/hexagon =>
backend/hexagon/runtime}/rpc/simulator/rpc_server.cc (99%)
rename src/{runtime/hexagon =>
backend/hexagon/runtime}/rpc/simulator/session.cc (99%)
rename src/{target/metal => backend/metal/codegen}/codegen_metal.cc (95%)
rename src/{target/metal => backend/metal/codegen}/codegen_metal.h (98%)
rename src/{target/metal => backend/metal/codegen}/intrin_rule_metal.cc (91%)
rename src/{target/metal => backend/metal/codegen}/metal_fallback_module.cc
(99%)
rename src/{target/metal => backend/metal/codegen}/metal_fallback_module.h
(98%)
create mode 100644 src/backend/metal/codegen/register.cc
copy tests/cpp/topi_ewise_test.cc => src/backend/metal/op/register.cc (71%)
create mode 100644 src/backend/metal/op/target_builtin.cc
rename src/{runtime/metal => backend/metal/runtime}/metal_common.h (99%)
rename src/{runtime/metal => backend/metal/runtime}/metal_device_api.mm (100%)
rename src/{runtime/metal => backend/metal/runtime}/metal_module.mm (98%)
rename src/{target/opencl => backend/opencl/codegen}/codegen_opencl.cc (98%)
rename src/{target/opencl => backend/opencl/codegen}/codegen_opencl.h (98%)
rename src/{target/opencl => backend/opencl/codegen}/intrin_rule_opencl.cc
(93%)
rename src/{target/opencl => backend/opencl/codegen}/opencl_fallback_module.cc
(97%)
rename src/{target/opencl => backend/opencl/codegen}/opencl_fallback_module.h
(98%)
create mode 100644 src/backend/opencl/codegen/register.cc
rename src/{runtime/opencl => backend/opencl/runtime}/opencl_common.h (99%)
rename src/{runtime/opencl => backend/opencl/runtime}/opencl_device_api.cc
(99%)
rename src/{runtime/opencl => backend/opencl/runtime}/opencl_module.cc (97%)
rename src/{runtime/opencl => backend/opencl/runtime}/opencl_wrapper/README.md
(100%)
rename src/{runtime/opencl =>
backend/opencl/runtime}/opencl_wrapper/opencl_wrapper.cc (100%)
rename src/{runtime/opencl => backend/opencl/runtime}/source_utils.h (100%)
rename src/{runtime/opencl => backend/opencl/runtime}/texture.h (100%)
rename src/{target/rocm => backend/rocm/codegen}/llvm/codegen_amdgpu.cc (97%)
rename src/{target/rocm => backend/rocm/codegen}/llvm/intrin_rule_rocm.cc (97%)
create mode 100644 src/backend/rocm/codegen/register.cc
rename src/{target/rocm => backend/rocm/codegen}/rocm_fallback_module.cc (99%)
rename src/{target/rocm => backend/rocm/codegen}/rocm_fallback_module.h (98%)
rename src/{runtime/rocm => backend/rocm/runtime}/rocm_common.h (98%)
rename src/{runtime/rocm => backend/rocm/runtime}/rocm_device_api.cc (100%)
rename src/{runtime/rocm => backend/rocm/runtime}/rocm_module.cc (98%)
rename src/{target/source => backend/trn/codegen}/codegen_trn.cc (91%)
rename src/{target/source => backend/trn/codegen}/codegen_trn.h (98%)
copy src/{tirx/transform/skip_assert.cc => backend/trn/codegen/register.cc}
(54%)
copy tests/cpp/topi_ewise_test.cc => src/backend/trn/op/register.cc (71%)
rename src/{tirx/op/target_builtin/trn.cc => backend/trn/op/target_builtin.cc}
(86%)
copy src/{tirx/transform/lower_tirx_cleanup.cc =>
backend/trn/transform/lower_trainium_layout.cc} (83%)
rename src/{target/vulkan => backend/vulkan/codegen}/build_vulkan.cc (88%)
rename src/{target/vulkan => backend/vulkan/codegen}/codegen_spirv.cc (98%)
rename src/{target/vulkan => backend/vulkan/codegen}/codegen_spirv.h (98%)
rename src/{target/vulkan => backend/vulkan/codegen}/intrin_rule_spirv.cc (92%)
rename src/{target/vulkan => backend/vulkan/codegen}/ir_builder.cc (100%)
rename src/{target/vulkan => backend/vulkan/codegen}/ir_builder.h (100%)
create mode 100644 src/backend/vulkan/codegen/register.cc
rename src/{target/vulkan => backend/vulkan/codegen}/spirv_support.cc (100%)
rename src/{target/vulkan => backend/vulkan/codegen}/spirv_support.h (100%)
rename src/{target/vulkan => backend/vulkan/codegen}/spirv_utils.cc (98%)
rename src/{target/vulkan => backend/vulkan/codegen}/spirv_utils.h (97%)
rename src/{target/vulkan => backend/vulkan/codegen}/vulkan_fallback_module.cc
(95%)
rename src/{target/vulkan => backend/vulkan/codegen}/vulkan_fallback_module.h
(98%)
rename src/{runtime/vulkan => backend/vulkan/runtime}/README.md (100%)
rename src/{runtime/vulkan => backend/vulkan/runtime}/spirv_shader.h (100%)
rename src/{runtime/vulkan => backend/vulkan/runtime}/thread_map.h (100%)
rename src/{runtime/vulkan => backend/vulkan/runtime}/vulkan_amdrgp.cc (100%)
rename src/{runtime/vulkan => backend/vulkan/runtime}/vulkan_amdrgp.h (100%)
rename src/{runtime/vulkan => backend/vulkan/runtime}/vulkan_buffer.cc (100%)
rename src/{runtime/vulkan => backend/vulkan/runtime}/vulkan_buffer.h (100%)
rename src/{runtime/vulkan => backend/vulkan/runtime}/vulkan_common.cc (100%)
rename src/{runtime/vulkan => backend/vulkan/runtime}/vulkan_common.h (100%)
rename src/{runtime/vulkan => backend/vulkan/runtime}/vulkan_device.cc (99%)
rename src/{runtime/vulkan => backend/vulkan/runtime}/vulkan_device.h (100%)
rename src/{runtime/vulkan => backend/vulkan/runtime}/vulkan_device_api.cc
(100%)
rename src/{runtime/vulkan => backend/vulkan/runtime}/vulkan_device_api.h (99%)
rename src/{runtime/vulkan => backend/vulkan/runtime}/vulkan_instance.cc (99%)
rename src/{runtime/vulkan => backend/vulkan/runtime}/vulkan_instance.h (100%)
rename src/{runtime/vulkan => backend/vulkan/runtime}/vulkan_module.cc (92%)
rename src/{runtime/vulkan => backend/vulkan/runtime}/vulkan_stream.cc (99%)
rename src/{runtime/vulkan => backend/vulkan/runtime}/vulkan_stream.h (100%)
rename src/{runtime/vulkan => backend/vulkan/runtime}/vulkan_wrapped_func.cc
(99%)
rename src/{runtime/vulkan => backend/vulkan/runtime}/vulkan_wrapped_func.h
(97%)
rename src/{target/webgpu => backend/webgpu/codegen}/codegen_webgpu.cc (98%)
rename src/{target/webgpu => backend/webgpu/codegen}/codegen_webgpu.h (98%)
rename src/{target/webgpu => backend/webgpu/codegen}/intrin_rule_webgpu.cc
(92%)
create mode 100644 src/backend/webgpu/codegen/register.cc
rename src/{target/webgpu => backend/webgpu/codegen}/webgpu_fallback_module.cc
(99%)
rename src/{target/webgpu => backend/webgpu/codegen}/webgpu_fallback_module.h
(98%)
