This is an automated email from the ASF dual-hosted git repository.
junrushao pushed a change to branch unity-staging
in repository https://gitbox.apache.org/repos/asf/tvm.git
from 525f5aee5f [Unity][TIR] Allow symbolic bounds in IndexMap analysis
(#15262)
add 399c5eaafe [Unity][Training] Registering te gradient (#15231)
add 6297535440 [MERGE] Hotfix layout transform related change after last
merge from main
add b00ed76318 [Unity][Dlight] Minor performance improvement for gemm and
gemv (#15278)
add d05c7e3844 [Unity][BYOC] `PrimValue` handling in `FuseOpByPattern` for
BYOC (#15217)
add ab8ac58241 [Unity][CUTLASS] Offload RMS norm (#15288)
add 6ab2e960fa [Unity] [DistIR] Introducing DistIR (#15289)
add fae8308813 [Dlight] Enhance fallback schedule with DecomposeReduction
(#15302)
add 5dcc25a15c [Unity][Relax][Transform] Do not remove MatchCast for
RemoveAllUnused (#15290)
add 1e45ca4ce3 [Unity] [Relax] [ONNX frontend] [op] Add support for Trilu
operator (#15299)
add 804ce09195 [Unity] Update CUTLASS Attention to incorprate upstream
change (#15309)
add f63fcba404 [Unity] Fix FuseOpsByPattern when a subgraph can be matched
by multiple residual patterns (#15308)
add 0924abad15 [Unity][Op] Add leaky relu operator (#15296)
add fd5a41f6f0 [Unity][Dlight] Add schedule rule for decode transpose
(#15304)
new b47b2695fb [Unity] CUDA Graph update (#15320)
new 98ef29a852 [Unity][Training] Enhance gradient system (#15230)
new aa28859340 [Unity][Dlight] Fix decode-GeMV rule when spatial-inner
without broadcasting (#15330)
new 6294aada46 [Unity][Dlight] Rule matmul avoiding blockIdx.z (#15333)
new 231653cca5 [Unity] fp16 A x int B GEMM update - support int8, more
bias shape (#15318)
new 959b7e5e09 [Unity][Dlight] Fix DecodeGeMV rule for spatial-inner with
grouping (#15340)
add 0413ce3138 [BugFix] Fix function to read all file (#15225)
add 30d684216c [TIR] Call TVMBackendFreeWorkspace inside LetStmt (#15253)
add 3c23865559 [Testing] Return BenchmarkResult in local_run and rpc_run
(#15277)
add dc7125b31e [Hexagon] Propagate QNN Concat Quantization Params to
Inputs (#15258)
add 9f8fe3c503 [topi] Add `arm_cpu` specific pooling schedules (#14855)
add e4a120955b [RELAY] Fix bug in MergeCompilerRegions pass (#15211)
add 81d7f79f03 Revert "[topi] Add `arm_cpu` specific pooling schedules"
(#15286)
add 592b3583dc [Exec] Add a script to test GPU memory bandwidth (#15287)
add fddbec7079 [TIR] Implement TIR macros (#15260)
add 33232deefb [FRONTEND][TFLITE][BugFix] Fix variable typo in batchmatmul
converting func (#15259)
add 1234f88b60 [BugFix][Relay][GraphExecutor] Fix set_input_zero_copy()
precision bug (#15291)
add 02ffc91396 [RPC] Fix socket bind errno on corner case (#15292)
add fba10d7021 [Docker] tensorflow_aarch64 package upgrade (#15293)
add a0e7d3e0ae [COMMUNITY] Qingchao Shen -> Reviewer (#15307)
add b6502f4e27 Fix keras version problem (#15265)
add 7890cca929 [JVM] Fix the Maven pom.xml for OS X arm64 tvm4j build
(#15321)
add 9af8efcd2d [Fix][TIR] LowerThreadAllreduce with correct thread mask
(#15323)
add 38bc953516 [Package] Remove cutlass media/docs inside
cutlass_fpA_intB_gemm (#15328)
add e25b1ba70a [TIR] ThreadAllreduce warp-level primitive support with
multi-warp (#15327)
add a4b863a2ff [Misc][Release] Extend PR tags and Format PR hyper-links in
release report (#15298)
add 7ad71e622a [Docker] Update ci-cortexm docker image (#15310)
add 6c63e0db53 [ETHOSU][MicroNPU][Pass] Add a pass to replicate pads
(#14909)
add c4f10cd5e9 [Runtime] Device API to query L2 cache size (#15332)
add e2d6511161 [Bugfix][Frontend][Keras]Fix a corner case bug in softmax
converter of keras frontend (#15337)
add c0946e19cd [Runtime] Flush L2 cache in time eval (#15305)
add 4b183daa97 [skipci] Fix typo in docs/arch/index.rst (#15312)
add a13b56a945 [OP] Add `rms_norm` into TOPI (#15326)
add d81e8809b8 [AOT] Avoid call_extern() with incorrect argument count
(#15301)
add 2eca9f0270 [TIR] Return error code from kernels in SplitHostDevice
(#15241)
new d8f1ac4e87 Merge remote-tracking branch 'apache-upstream/main' into
unity
The 7 revisions listed above as "new" are entirely new to this
repository and will be described in separate emails. The revisions
listed as "add" were already present in the repository and have only
been added to this reference.
Summary of changes:
3rdparty/cutlass | 2 +-
3rdparty/cutlass_fpA_intB_gemm | 2 +-
CMakeLists.txt | 3 +
CONTRIBUTORS.md | 1 +
.../template_project/microtvm_api_server.py | 9 +-
ci/jenkins/docker-images.ini | 2 +-
.../install/ubuntu_install_tensorflow_aarch64.sh | 2 +-
docs/arch/index.rst | 4 +-
include/tvm/ir/global_info.h | 2 +-
include/tvm/relax/attrs/distributed.h | 48 +
include/tvm/relax/attrs/nn.h | 9 +
include/tvm/relax/attrs/op.h | 49 +
include/tvm/relax/distributed/axis_group_graph.h | 305 ++++++
include/tvm/relax/distributed/global_info.h | 92 ++
include/tvm/relax/distributed/struct_info.h | 192 ++++
include/tvm/relax/distributed/transform.h | 56 +
include/tvm/relax/struct_info.h | 2 +-
include/tvm/relax/struct_info_functor.h | 6 +
include/tvm/runtime/device_api.h | 3 +-
include/tvm/runtime/profiling.h | 4 +-
include/tvm/topi/nn/rms_norm.h | 96 ++
jvm/native/osx-x86_64/pom.xml | 2 +-
jvm/pom.xml | 2 +-
python/setup.py | 8 +
python/tvm/_ffi/runtime_ctypes.py | 18 +
python/tvm/contrib/cutlass/attention_operation.py | 8 +-
python/tvm/contrib/cutlass/build.py | 34 +-
python/tvm/contrib/cutlass/conv2d_operation.py | 7 +-
python/tvm/contrib/cutlass/gemm_operation.py | 73 +-
python/tvm/contrib/cutlass/gen_tensor_op.py | 37 +-
python/tvm/contrib/cutlass/layer_norm_operation.py | 10 +-
...yer_norm_operation.py => rms_norm_operation.py} | 19 +-
python/tvm/contrib/hexagon/transform.py | 105 +-
python/tvm/dlight/base/__init__.py | 8 +-
python/tvm/dlight/base/analysis.py | 46 +-
python/tvm/dlight/gpu/__init__.py | 1 +
python/tvm/dlight/gpu/decode_gemv.py | 63 +-
python/tvm/dlight/gpu/fallback.py | 10 +-
python/tvm/dlight/gpu/matmul.py | 15 +-
python/tvm/dlight/gpu/transpose.py | 129 +++
python/tvm/exec/gpu_memory_bandwidth.py | 192 ++++
python/tvm/relax/__init__.py | 13 +-
python/tvm/relax/analysis/analysis.py | 11 +-
python/tvm/relax/backend/contrib/cutlass.py | 60 +-
python/tvm/relax/backend/patterns.py | 22 +-
python/tvm/relax/block_builder.py | 59 +-
.../relax => relax/distributed}/__init__.py | 11 +-
.../__init__.py => relax/distributed/_ffi_api.py} | 8 +-
python/tvm/relax/distributed/global_info.py | 70 ++
python/tvm/relax/distributed/struct_info.py | 141 +++
.../distributed/transform}/__init__.py | 8 +-
.../distributed/transform/_ffi_api.py} | 9 +-
.../tvm/relax/distributed/transform/transform.py | 23 +-
python/tvm/relax/frontend/onnx/onnx_frontend.py | 18 +
python/tvm/relax/frontend/torch/fx_translator.py | 13 +
python/tvm/relax/op/__init__.py | 3 +-
python/tvm/relax/op/_op_gradient.py | 35 +-
python/tvm/relax/op/base.py | 62 +-
.../relax => relax/op/distributed}/__init__.py | 8 +-
.../op/distributed/_ffi_api.py} | 8 +-
python/tvm/relax/op/distributed/distributed.py | 61 ++
python/tvm/relax/op/grad/grad.py | 53 +
python/tvm/relax/op/manipulate.py | 4 +-
python/tvm/relax/op/nn/nn.py | 23 +
python/tvm/relax/op/op_attrs.py | 5 +
python/tvm/relax/testing/nn.py | 123 ++-
python/tvm/relax/training/utils.py | 56 +-
python/tvm/relax/transform/legalize_ops/grad.py | 15 +-
.../tvm/relax/transform/legalize_ops/manipulate.py | 5 +-
python/tvm/relax/transform/legalize_ops/nn.py | 5 +
python/tvm/relax/transform/transform.py | 4 +
python/tvm/relay/backend/contrib/ethosu/codegen.py | 89 +-
python/tvm/relay/frontend/keras.py | 21 +-
python/tvm/relay/frontend/tflite.py | 12 +-
python/tvm/relay/op/contrib/ethosu.py | 4 +-
python/tvm/runtime/module.py | 5 +
python/tvm/script/ir_builder/relax/__init__.py | 1 +
.../ir_builder/relax/{ => distributed}/__init__.py | 3 +-
.../relax/{__init__.py => distributed/_ffi_api.py} | 10 +-
.../tvm/script/ir_builder/relax/distributed/ir.py | 160 +++
python/tvm/script/ir_builder/relax/ir.py | 2 +
python/tvm/script/parser/_core.py | 2 +-
python/tvm/script/parser/core/entry.py | 35 +-
python/tvm/script/parser/ir/__init__.py | 10 +-
python/tvm/script/parser/relax/__init__.py | 27 +-
python/tvm/script/parser/relax/dist.py | 95 ++
python/tvm/script/parser/tir/__init__.py | 4 +-
python/tvm/script/parser/tir/entry.py | 99 +-
python/tvm/script/parser/tir/parser.py | 60 +-
python/tvm/target/target.py | 4 +
python/tvm/testing/runner.py | 12 +-
python/tvm/tir/op.py | 2 +-
python/tvm/topi/nn/__init__.py | 1 +
python/tvm/topi/nn/rms_norm.py | 46 +
python/tvm/topi/testing/__init__.py | 1 +
python/tvm/topi/testing/rms_norm_python.py | 51 +
src/relax/analysis/struct_info_analysis.cc | 25 +
.../backend/contrib/codegen_json/codegen_json.h | 2 +-
src/relax/distributed/axis_group_graph.cc | 478 +++++++++
src/relax/distributed/global_info.cc | 70 ++
src/relax/distributed/struct_info.cc | 141 +++
.../distributed/transform/propagate_sharding.cc | 512 ++++++++++
src/relax/ir/binding_rewrite.cc | 2 +-
src/relax/ir/block_builder.cc | 18 +
src/relax/ir/struct_info_functor.cc | 10 +
src/relax/op/distributed/binary.cc | 74 ++
src/relax/op/distributed/binary.h | 92 ++
src/relax/op/distributed/distributed.cc | 89 ++
src/relax/op/distributed/distributed.h | 57 ++
src/relax/op/distributed/linear_algebra.cc | 106 ++
src/relax/op/distributed/linear_algebra.h | 39 +
src/relax/op/distributed/manipulate.cc | 145 +++
src/relax/op/distributed/manipulate.h | 41 +
src/relax/op/distributed/nn.cc | 66 ++
src/relax/op/distributed/nn.h | 40 +
src/relax/op/distributed/op.cc | 42 +
src/relax/op/distributed/statistical.cc | 90 ++
src/relax/op/distributed/statistical.h | 44 +
src/relax/op/distributed/unary.cc | 64 ++
src/relax/op/distributed/unary.h | 78 ++
src/relax/op/distributed/utils.cc | 85 ++
src/relax/op/distributed/utils.h | 67 ++
src/relax/op/nn/nn.cc | 20 +
src/relax/op/nn/nn.h | 3 +
src/relax/op/op.cc | 51 +
src/relax/op/op_common.cc | 11 +
src/relax/op/op_common.h | 23 +
src/relax/op/tensor/grad.cc | 47 +-
src/relax/transform/fuse_ops.cc | 28 +-
src/relax/transform/gradient.cc | 591 ++++++++---
src/relax/transform/rewrite_cuda_graph.cc | 20 +-
src/relay/backend/aot_executor_codegen.cc | 38 +-
src/relay/transforms/merge_compiler_regions.cc | 36 +-
src/runtime/contrib/cutlass/weight_preprocess.cc | 11 +-
src/runtime/crt/common/crt_runtime_api.c | 5 +-
src/runtime/cuda/cuda_device_api.cc | 10 +
.../graph_executor/debug/graph_executor_debug.cc | 2 +-
src/runtime/graph_executor/graph_executor.cc | 25 +-
src/runtime/graph_executor/graph_executor.h | 6 +-
src/runtime/metal/metal_device_api.mm | 2 +
src/runtime/opencl/opencl_device_api.cc | 7 +
src/runtime/profiling.cc | 16 +-
src/runtime/relax_vm/cuda/cuda_graph_builtin.cc | 47 +-
src/runtime/rocm/rocm_device_api.cc | 5 +
src/runtime/rpc/rpc_module.cc | 28 +-
src/runtime/vulkan/vulkan_device_api.cc | 3 +
src/script/ir_builder/relax/distributed.cc | 61 ++
src/script/printer/ir/distributed.cc | 42 +
src/script/printer/ir/ir.cc | 11 +
src/script/printer/ir/utils.h | 7 +-
src/script/printer/relax/call.cc | 39 +-
src/script/printer/relax/distributed.cc | 132 +++
src/script/printer/relax/expr.cc | 7 +
src/script/printer/relax/tir.cc | 9 +
src/script/printer/tir/ir.cc | 5 +-
src/support/socket.h | 16 +-
src/te/operation/cross_thread_reduction.cc | 13 +-
src/tir/transforms/lower_device_kernel_launch.cc | 41 +-
src/tir/transforms/lower_thread_allreduce.cc | 337 +++---
src/tir/transforms/lower_tvm_builtin.cc | 54 +-
src/tir/transforms/split_host_device.cc | 33 +-
src/topi/nn.cc | 6 +
tests/python/contrib/test_ethosu/test_codegen.py | 63 ++
tests/python/contrib/test_ethosu/test_legalize.py | 132 ++-
.../test_hexagon/test_relay_simplify_qnn_concat.py | 101 ++
tests/python/dlight/test_gpu_decode_gemv.py | 244 ++++-
tests/python/dlight/test_gpu_fallback.py | 47 +-
tests/python/dlight/test_gpu_matmul.py | 215 ++--
tests/python/dlight/test_gpu_transpose.py | 189 ++++
tests/python/frontend/keras/test_forward.py | 7 +
tests/python/frontend/tflite/test_forward.py | 18 +
.../distributed/test_distributed_dtensor_sinfo.py | 93 ++
...est_distributed_transform_propagate_sharding.py | 1077 ++++++++++++++++++++
.../test_distributed_tvmscript_parser.py | 193 ++++
.../test_distributed_tvmscript_printer.py | 159 +++
tests/python/relax/test_analysis.py | 23 +
.../test_analysis_suggest_layout_transforms.py | 1 +
tests/python/relax/test_codegen_cutlass.py | 326 +++++-
tests/python/relax/test_frontend_from_fx.py | 37 +
tests/python/relax/test_frontend_onnx.py | 8 +
tests/python/relax/test_op_grad.py | 45 +
tests/python/relax/test_op_gradient_numeric.py | 24 +-
tests/python/relax/test_op_misc.py | 19 +
tests/python/relax/test_op_nn.py | 6 +
tests/python/relax/test_training_setup_trainer.py | 25 +-
.../relax/test_transform_fuse_ops_by_pattern.py | 150 ++-
tests/python/relax/test_transform_gradient.py | 949 +++++++++--------
.../relax/test_transform_gradient_checkpoint.py | 689 +++++++++++++
.../relax/test_transform_gradient_te_register.py | 384 +++++++
.../relax/test_transform_lazy_transform_params.py | 3 +
.../python/relax/test_transform_legalize_ops_nn.py | 73 ++
.../relax/test_transform_rewrite_cuda_graph.py | 322 ++++++
tests/python/relax/test_tvmscript_parser.py | 28 +
tests/python/relax/test_tvmscript_printer_relax.py | 21 +
.../relay/test_pass_merge_compiler_regions.py | 62 ++
tests/python/topi/python/test_topi_rms_norm.py | 68 ++
tests/python/unittest/test_set_input_zero_copy.py | 137 +++
.../test_tir_transform_lower_thread_all_reduce.py | 451 ++++++++
.../test_tir_transform_lower_tvm_builtin.py | 23 +-
.../test_tir_transform_split_host_device.py | 38 +
tests/python/unittest/test_tvmscript_parser_tir.py | 107 ++
tests/scripts/release/README.md | 13 +-
tests/scripts/release/make_notes.py | 33 +-
web/emcc/tvmjs_support.cc | 2 +-
204 files changed, 12984 insertions(+), 1314 deletions(-)
create mode 100644 include/tvm/relax/attrs/distributed.h
create mode 100644 include/tvm/relax/attrs/op.h
create mode 100644 include/tvm/relax/distributed/axis_group_graph.h
create mode 100644 include/tvm/relax/distributed/global_info.h
create mode 100644 include/tvm/relax/distributed/struct_info.h
create mode 100644 include/tvm/relax/distributed/transform.h
create mode 100644 include/tvm/topi/nn/rms_norm.h
copy python/tvm/contrib/cutlass/{layer_norm_operation.py =>
rms_norm_operation.py} (72%)
create mode 100644 python/tvm/dlight/gpu/transpose.py
create mode 100644 python/tvm/exec/gpu_memory_bandwidth.py
copy python/tvm/{script/ir_builder/relax => relax/distributed}/__init__.py
(79%)
copy python/tvm/{script/ir_builder/relax/__init__.py =>
relax/distributed/_ffi_api.py} (82%)
create mode 100644 python/tvm/relax/distributed/global_info.py
create mode 100644 python/tvm/relax/distributed/struct_info.py
copy python/tvm/{script/ir_builder/relax =>
relax/distributed/transform}/__init__.py (82%)
copy python/tvm/{script/ir_builder/relax/__init__.py =>
relax/distributed/transform/_ffi_api.py} (80%)
copy docker/install/ubuntu_install_tensorflow_aarch64.sh =>
python/tvm/relax/distributed/transform/transform.py (68%)
mode change 100755 => 100644
copy python/tvm/{script/ir_builder/relax => relax/op/distributed}/__init__.py
(82%)
copy python/tvm/{script/ir_builder/relax/__init__.py =>
relax/op/distributed/_ffi_api.py} (82%)
create mode 100644 python/tvm/relax/op/distributed/distributed.py
copy python/tvm/script/ir_builder/relax/{ => distributed}/__init__.py (93%)
copy python/tvm/script/ir_builder/relax/{__init__.py =>
distributed/_ffi_api.py} (80%)
create mode 100644 python/tvm/script/ir_builder/relax/distributed/ir.py
create mode 100644 python/tvm/script/parser/relax/dist.py
create mode 100644 python/tvm/topi/nn/rms_norm.py
create mode 100644 python/tvm/topi/testing/rms_norm_python.py
create mode 100644 src/relax/distributed/axis_group_graph.cc
create mode 100644 src/relax/distributed/global_info.cc
create mode 100644 src/relax/distributed/struct_info.cc
create mode 100644 src/relax/distributed/transform/propagate_sharding.cc
create mode 100644 src/relax/op/distributed/binary.cc
create mode 100644 src/relax/op/distributed/binary.h
create mode 100644 src/relax/op/distributed/distributed.cc
create mode 100644 src/relax/op/distributed/distributed.h
create mode 100644 src/relax/op/distributed/linear_algebra.cc
create mode 100644 src/relax/op/distributed/linear_algebra.h
create mode 100644 src/relax/op/distributed/manipulate.cc
create mode 100644 src/relax/op/distributed/manipulate.h
create mode 100644 src/relax/op/distributed/nn.cc
create mode 100644 src/relax/op/distributed/nn.h
create mode 100644 src/relax/op/distributed/op.cc
create mode 100644 src/relax/op/distributed/statistical.cc
create mode 100644 src/relax/op/distributed/statistical.h
create mode 100644 src/relax/op/distributed/unary.cc
create mode 100644 src/relax/op/distributed/unary.h
create mode 100644 src/relax/op/distributed/utils.cc
create mode 100644 src/relax/op/distributed/utils.h
create mode 100644 src/script/ir_builder/relax/distributed.cc
create mode 100644 src/script/printer/ir/distributed.cc
create mode 100644 src/script/printer/relax/distributed.cc
create mode 100644
tests/python/contrib/test_hexagon/test_relay_simplify_qnn_concat.py
create mode 100644 tests/python/dlight/test_gpu_transpose.py
create mode 100644
tests/python/relax/distributed/test_distributed_dtensor_sinfo.py
create mode 100644
tests/python/relax/distributed/test_distributed_transform_propagate_sharding.py
create mode 100644
tests/python/relax/distributed/test_distributed_tvmscript_parser.py
create mode 100644
tests/python/relax/distributed/test_distributed_tvmscript_printer.py
create mode 100644 tests/python/relax/test_transform_gradient_checkpoint.py
create mode 100644 tests/python/relax/test_transform_gradient_te_register.py
create mode 100644 tests/python/topi/python/test_topi_rms_norm.py
create mode 100644 tests/python/unittest/test_set_input_zero_copy.py
