This is an automated email from the ASF dual-hosted git repository.
tqchen pushed a change to branch unity-staging
in repository https://gitbox.apache.org/repos/asf/tvm.git
from 153f70aa60 [MERGE] Merge main into unity 2023-06-01
add 2f13f2d3bf [Unity] Fix importing tvm.contrib.cutlass (#15010)
add 6d8fe1c0e3 [Unity] Fix cutlass BYOC after merge (#15012)
add 1a4e08e441 [Unity][FIX] add init file to `relax.backend.contrib`
(#15023)
add 3dd9345024 [Unity] Allow name_hint in additional locations (#15027)
add 7dd76cca03 [Unity] Cutlass attention with dynamic sequence length
(#15028)
add 91b6cee634 [Unity][BYOC] Cache cuBlasLt handle with thread entry
(#15030)
add 0328e98b94 [Unity][NN] Allow nn.Placeholder/Parameter prior to
BlockBuilder (#15025)
add 7c15b55aaa [Unity] Add hexp for compute capacity <= 5.2 (#15070)
add e9ddd47a80 [Unity] Add an API to create multiple kv caches with single
allocation (#15064)
add 0406b9af22 [Unity][Frontend] Add relax onnx importer and tests (#14999)
add bbe5d6b18c [Unity] Added bounds checking on TupleGetItem index (#15024)
add f61015c20c [Unity] Optimize SampleTopPFromProb (#15072)
add dc1e580e61 [Unity][Relax] Add masked_fill operator (#15077)
add c1aab7d185 [Unity][Relax] Add bitwise and logical ops (AND, NOT, OR,
XOR) (#15075)
add 208dd71f18 [Bugfix][CUDA] Fix codegen for hexp for sm >= 52 (#15079)
add 4347b30ffe [Unity] Fix FewShotTuning Failure When Missing Global
Symbol (#15097)
add 00126b0484 [IR,TE,TIR] Use f-strings for string formatting, NFC
(#14990)
add 521465e626 [Minor] Fix compilation warnings (#15007)
add af5a6c48e3 [CodegenC] Updated unit test for sorted CodegenC output
(#14949)
add 2618091ae7 [Hotfix] Remove `LOG(INFO)` from unsupported dtype
legalization pass (#15015)
add 0af9ff90e0 [TIR] Restrict tir.transform.LowerTVMBuiltin to host
functions (#14944)
add 5b9e9bd675 [TIR] Keep trivial LetStmt in tir.Simplify when used in
buffer decl (#14951)
add 9877db5ff3 [TIR] Handle callees on same target, different codegen
(#14988)
add 429f601284 [TIR] Enhance TVMScript Buffer Slice Access (#14693)
add 6bcc1b44f5 [TOPI] Fix dynamic dimensions support for Dense on TOPI
side (#15018)
add efb09a5953 [ONNX] Fix bug in scatter_elements (#15017)
add 80079b6493 Fix type parse error about AdaptiveMaxPool (#15016)
add 1366f2e250 [OpenCLML] Reactor and introduce on chip memory and memory
planner (#14922)
add 7b367da281 [Frontend][Paddle] [PaddlePaddle Hackathon 4] add convert
support for tanhshrink/pool3d/set_value ops for paddle frontend (#14973)
add 3dba48bfdc [Frontend][Paddle] [PaddlePaddle Hackathon 4]add attribute
support for gaussian_random/softplus/Conv3d/Conv2d (#14801)
add a74a2e2d2e [Build] Fix missing virtual destructor in SIBuilder (#15022)
add b59e866c13 [BYOC] Hide internal cutlass symbols (#15029)
add d2d9dc7c04 [TIR] Preserve object equality in
Buffer::GetFlattenedBuffer (#15040)
add 26907f9d99 [Vulkan] Allow DeclBuffer in CodeGenSPIRV (#15035)
add 3d72d4ba0e [StackVM] Updated CodeGenStackVM to handle DeclBuffer
(#15036)
add bf91083918 [TIR] Handle DeclBuffer in StorageFlatten's input (#15042)
add 2d2b727339 [Docker] Update ci_cpu docker image to one containing polly
package f… (#15031)
add 20aeebdec0 [#14992][DOC][TUTORIAL] Fix typo for the 'Making your
Hardware Accelerator TVM-ready with UMA' (#15055)
add 56255f8dc5 [TIR] Move SplitHostDevice to before MakePackedAPI (#14986)
add fa4d73ac51 [TIR] Handle DeclBuffer in IRSubstitute (#15054)
add e5c1a24bda [TIR] Refactor BlockScope outside schedule (#15034)
add f3fde8121b [TIR] [Schedule] Fix decompose_padding bug with dtypes
(#15050)
add e6f0a15712 [Bugfix][CI] Fix CacheControl version constraint violation
(#15056)
add 7dc1d6a99f [CI] Make Graviton3 default AArch64 job runner node (#14983)
add 52c3d84040 [TIR] Handle DeclBuffer in StorageRewrite (#15051)
add 9e265ff5f0 [Bugfix][Relay][keras] Fix the bug about the attribute
'output_padding' in Deconv (#15060)
add f172f6c70d Revert "[CI] Make Graviton3 default AArch64 job runner node
(#14983)" (#15065)
add df5b180ebd [TOPI][Adreno] Fix problem with ceil_log2 (#15061)
add 997adcaaea [CI] Clone submodule recursively (#15062)
add 5fca2b2fd3 [TIR][Analysis] Handle DeclBuffer in FlopEstimator (#15052)
add eea6268c92 [TIR] Handle DeclBuffer in
Inline/ComputeAt/ReverseComputeAt (#15038)
add dee3c2ab28 [TIR] Handle DeclBuffer in LowerCustomDatatypes (#15041)
add d8e5812db7 [BYOC] Add GEMM kernel from FasterTransformer as submodule
(#15046)
add 21361a63aa [CMSIS-NN] Update CMSIS-NN release to v4.1.0 (#15059)
add 6c6ad6c43a [AOT] Remove duplication in tvm.testing.aot.compile_models
(#15032)
add 70532b87d1 [COMMUNITY] Aleksei-grovety -> Reviewer (#15086)
add 68ac909713 [CodeGenC] Use PrimFuncNode::ret_type in function signature
(#15073)
add 02dc191e86 [ARITH] Improve arith simplify to handle symbolic reshape
pattern (#15081)
new 9eafe7c022 [MERGE] Merge main into unity 2023-06-13
The 1 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:
.gitmodules | 3 +
3rdparty/cutlass_fpA_intB_gemm | 1 +
CMakeLists.txt | 8 +-
CONTRIBUTORS.md | 1 +
apps/cpp_clml/clml_runner.cc | 36 +-
apps/cpp_clml/clml_runner.h | 2 +-
apps/cpp_clml/scripts/clml_codegen.py | 2 +-
ci/jenkins/docker-images.ini | 2 +-
ci/jenkins/generated/arm_jenkinsfile.groovy | 4 +-
ci/jenkins/generated/cortexm_jenkinsfile.groovy | 4 +-
ci/jenkins/generated/cpu_jenkinsfile.groovy | 4 +-
ci/jenkins/generated/docker_jenkinsfile.groovy | 4 +-
ci/jenkins/generated/gpu_jenkinsfile.groovy | 4 +-
ci/jenkins/generated/hexagon_jenkinsfile.groovy | 4 +-
ci/jenkins/generated/i386_jenkinsfile.groovy | 4 +-
ci/jenkins/generated/lint_jenkinsfile.groovy | 4 +-
.../generated/minimal_cross_isa_jenkinsfile.groovy | 4 +-
ci/jenkins/generated/minimal_jenkinsfile.groovy | 4 +-
ci/jenkins/generated/riscv_jenkinsfile.groovy | 4 +-
ci/jenkins/generated/wasm_jenkinsfile.groovy | 4 +-
ci/jenkins/templates/utils/Prepare.groovy.j2 | 2 +-
cmake/modules/OpenCL.cmake | 39 +-
cmake/modules/contrib/CUTLASS.cmake | 3 +
docker/install/ubuntu_install_cmsis.sh | 2 +-
.../python/bootstrap/lockfiles/constraints-3.7.txt | 6 +-
.../python/bootstrap/lockfiles/constraints-3.8.txt | 6 +-
gallery/tutorial/uma.py | 2 +-
include/tvm/relax/nested_msg.h | 2 -
include/tvm/tir/{schedule => }/block_scope.h | 18 +-
include/tvm/tir/schedule/state.h | 20 +-
include/tvm/tir/utils.h | 96 +
...lpack.txt => LICENSE.cutlass_fpA_intB_gemm.txt} | 5 +-
python/tvm/contrib/cutlass/attention_operation.py | 17 +-
python/tvm/contrib/cutlass/build.py | 1 +
python/tvm/contrib/cutlass/conv2d_operation.py | 3 -
python/tvm/contrib/cutlass/gemm_operation.py | 3 -
python/tvm/contrib/cutlass/gen_tensor_op.py | 37 +-
python/tvm/contrib/hexagon/pytest_plugin.py | 2 +-
python/tvm/ir/container.py | 4 +-
python/tvm/ir/expr.py | 6 +-
python/tvm/ir/json_compact.py | 4 +-
python/tvm/ir/supply.py | 7 +-
.../debugger => relax/backend/contrib}/__init__.py | 1 +
python/tvm/relax/backend/contrib/cutlass.py | 7 +-
python/tvm/relax/block_builder.py | 3 +-
python/tvm/relax/expr.py | 12 +-
.../tvm/relax/frontend/onnx}/__init__.py | 6 +-
python/tvm/relax/frontend/onnx/onnx_frontend.py | 2173 ++++++++++++++++++++
python/tvm/relax/op/__init__.py | 1 +
python/tvm/relax/op/binary.py | 102 +
python/tvm/relax/op/{ternary.py => mask.py} | 29 +-
python/tvm/relax/op/unary.py | 32 +
python/tvm/relax/testing/nn.py | 47 +-
python/tvm/relax/transform/legalize_ops/binary.py | 10 +
python/tvm/relax/transform/legalize_ops/unary.py | 2 +
python/tvm/relay/frontend/keras.py | 4 +
python/tvm/relay/frontend/onnx.py | 2 +-
python/tvm/relay/frontend/paddlepaddle.py | 189 +-
python/tvm/relay/frontend/pytorch.py | 4 +
python/tvm/relay/op/contrib/clml.py | 165 +-
python/tvm/script/ir_builder/relax/ir.py | 16 +
python/tvm/script/parser/core/evaluator.py | 38 +
python/tvm/script/parser/tir/operation.py | 92 +-
python/tvm/script/parser/tir/parser.py | 22 +
python/tvm/te/hybrid/module.py | 2 +-
python/tvm/te/hybrid/parser.py | 10 +-
python/tvm/te/hybrid/preprocessor.py | 4 +-
python/tvm/te/hybrid/utils.py | 6 +-
python/tvm/te/operation.py | 11 +-
python/tvm/te/schedule.py | 2 +-
python/tvm/te/tag.py | 2 +-
python/tvm/te/tensor.py | 4 +-
python/tvm/testing/aot.py | 83 +-
python/tvm/tir/{schedule => }/block_scope.py | 0
python/tvm/tir/buffer.py | 15 +-
python/tvm/tir/ir_builder.py | 2 +-
python/tvm/tir/schedule/__init__.py | 2 +-
python/tvm/tir/schedule/schedule.py | 170 +-
python/tvm/tir/schedule/state.py | 11 +-
python/tvm/tir/tensor_intrin/arm_cpu.py | 8 +-
python/tvm/tir/tensor_intrin/cuda.py | 17 +-
python/tvm/topi/math.py | 3 +
python/tvm/topi/nn/dense.py | 12 +-
python/tvm/utils/roofline/__init__.py | 12 +-
src/arith/iter_affine_map.cc | 11 +-
src/driver/driver_api.cc | 5 +-
src/ir/si_builder.cc | 1 +
src/relax/backend/vm/vm_shape_lower.cc | 6 +-
src/relax/ir/expr.cc | 11 +
src/relax/op/tensor/binary.cc | 12 +
src/relax/op/tensor/binary.h | 22 +
src/relax/op/tensor/unary.cc | 2 +
src/relax/op/tensor/unary.h | 6 +
src/relax/transform/few_shot_tuning.cc | 3 +-
src/relay/printer/relay_text_printer.cc | 2 +-
src/runtime/contrib/clml/clml_memory_planner.cc | 268 +++
.../runtime/contrib/clml/clml_memory_planner.h | 21 +-
src/runtime/contrib/clml/clml_runtime.cc | 1035 +++++-----
src/runtime/contrib/clml/clml_runtime.h | 185 ++
src/runtime/contrib/clml/clml_utils.cc | 257 +++
src/runtime/contrib/clml/clml_utils.h | 77 +
src/runtime/contrib/cublas/cublas_json_runtime.cc | 9 +-
src/runtime/contrib/cublas/cublas_utils.cc | 13 +
src/runtime/contrib/cublas/cublas_utils.h | 7 +
src/runtime/relax_vm/lm_support.cc | 139 +-
src/runtime/relax_vm/memory_manager.cc | 7 +
src/runtime/vulkan/vulkan_amdrgp.cc | 8 +-
src/runtime/vulkan/vulkan_buffer.cc | 2 +-
src/runtime/vulkan/vulkan_instance.cc | 4 +-
src/runtime/vulkan/vulkan_stream.cc | 6 +-
src/runtime/vulkan/vulkan_wrapped_func.cc | 16 +-
src/target/llvm/codegen_cpu.cc | 2 +-
src/target/source/codegen_c.cc | 10 +-
src/target/source/codegen_c.h | 7 +-
src/target/source/codegen_c_host.cc | 15 +-
src/target/source/codegen_c_host.h | 5 +-
src/target/source/codegen_cuda.cc | 2 +-
src/target/source/codegen_opencl.cc | 2 +-
src/target/source/codegen_vhls.cc | 2 +-
src/target/source/literal/cuda_half_t.h | 15 +-
src/target/spirv/codegen_spirv.cc | 2 +
src/target/spirv/codegen_spirv.h | 1 +
src/target/stackvm/codegen_stackvm.cc | 2 +
src/target/stackvm/codegen_stackvm.h | 1 +
src/target/target_kind.cc | 5 +-
src/tir/analysis/estimate_flops.cc | 1 +
src/tir/{schedule => ir}/block_scope.cc | 25 +-
src/tir/ir/buffer.cc | 17 +-
src/tir/ir/stmt_functor.cc | 26 +-
src/tir/schedule/analysis/analysis.cc | 2 +-
src/tir/schedule/analysis/verify.cc | 6 +-
src/tir/schedule/concrete_schedule.cc | 1 -
src/tir/schedule/primitive/blockize_tensorize.cc | 2 -
src/tir/schedule/primitive/cache_index.cc | 4 +-
src/tir/schedule/primitive/cache_read_write.cc | 14 +-
src/tir/schedule/primitive/decompose_padding.cc | 6 +-
src/tir/schedule/primitive/read_write_at.cc | 2 +-
src/tir/schedule/primitive/reduction.cc | 2 +-
src/tir/schedule/state.cc | 9 +-
src/tir/schedule/transform.cc | 28 +-
src/tir/schedule/utils.h | 71 +-
src/tir/transforms/ir_utils.cc | 5 +
src/tir/transforms/lower_custom_datatypes.cc | 5 +
src/tir/transforms/lower_device_kernel_launch.cc | 45 +-
src/tir/transforms/lower_tvm_builtin.cc | 8 +-
src/tir/transforms/make_packed_api.cc | 11 +-
src/tir/transforms/make_unpacked_api.cc | 10 +-
src/tir/transforms/simplify.cc | 88 +-
src/tir/transforms/split_host_device.cc | 8 -
src/tir/transforms/storage_flatten.cc | 9 +
src/tir/transforms/storage_rewrite.cc | 70 +-
src/tir/transforms/unsupported_dtype_legalize.cc | 1 -
tests/cpp-runtime/opencl/clml_memory_planner.cc | 439 ++++
tests/cpp/c_codegen_test.cc | 8 +-
tests/python/contrib/test_clml/infrastructure.py | 19 +
tests/python/contrib/test_clml/test_ops.py | 110 +-
.../contrib/test_ethosu/test_encode_constants.py | 4 +-
.../test_ethosu/test_tir_to_cs_translator.py | 8 +-
tests/python/frontend/keras/test_forward.py | 6 +
tests/python/frontend/paddlepaddle/test_forward.py | 183 +-
tests/python/frontend/pytorch/test_forward.py | 7 +
tests/python/relax/test_blockbuilder_core.py | 31 +
tests/python/relax/test_codegen_cutlass.py | 74 +-
tests/python/relax/test_frontend_onnx.py | 1619 +++++++++++++++
tests/python/relax/test_op_binary.py | 12 +
tests/python/relax/test_op_unary.py | 6 +
tests/python/relax/test_runtime_builtin.py | 28 +-
.../python/relax/test_transform_few_shot_tuning.py | 14 +-
.../python/unittest/test_arith_iter_affine_map.py | 41 +-
tests/python/unittest/test_custom_datatypes.py | 260 ++-
.../python/unittest/test_target_codegen_opencl.py | 27 +
.../unittest/test_target_codegen_vm_basic.py | 22 +-
.../python/unittest/test_target_codegen_vulkan.py | 19 +-
.../test_tir_analysis_estimate_tir_flops.py | 17 +
tests/python/unittest/test_tir_buffer.py | 23 +
.../unittest/test_tir_schedule_compute_at.py | 99 +
.../test_tir_schedule_decompose_padding.py | 41 +
.../unittest/test_tir_stmt_functor_substitute.py | 83 +
.../test_tir_transform_device_kernel_launch.py | 49 +
.../test_tir_transform_lower_tvm_builtin.py | 97 +-
.../unittest/test_tir_transform_make_packed_api.py | 40 +-
.../test_tir_transform_make_unpacked_api.py | 65 +-
.../python/unittest/test_tir_transform_simplify.py | 44 +
.../test_tir_transform_split_host_device.py | 14 +-
.../unittest/test_tir_transform_storage_flatten.py | 16 +
.../unittest/test_tir_transform_storage_rewrite.py | 66 +
tests/python/unittest/test_tvmscript_ops.py | 69 +
tests/scripts/task_build_adreno_bins.sh | 4 +-
188 files changed, 8672 insertions(+), 1395 deletions(-)
create mode 160000 3rdparty/cutlass_fpA_intB_gemm
rename include/tvm/tir/{schedule => }/block_scope.h (93%)
create mode 100644 include/tvm/tir/utils.h
copy licenses/{LICENSE.dlpack.txt => LICENSE.cutlass_fpA_intB_gemm.txt} (99%)
copy python/tvm/{contrib/debugger => relax/backend/contrib}/__init__.py (96%)
copy {apps/uma/_template => python/tvm/relax/frontend/onnx}/__init__.py (89%)
create mode 100644 python/tvm/relax/frontend/onnx/onnx_frontend.py
copy python/tvm/relax/op/{ternary.py => mask.py} (63%)
rename python/tvm/tir/{schedule => }/block_scope.py (100%)
create mode 100644 src/runtime/contrib/clml/clml_memory_planner.cc
copy include/tvm/runtime/contrib/libtorch_runtime.h =>
src/runtime/contrib/clml/clml_memory_planner.h (65%)
create mode 100644 src/runtime/contrib/clml/clml_runtime.h
create mode 100644 src/runtime/contrib/clml/clml_utils.cc
create mode 100644 src/runtime/contrib/clml/clml_utils.h
rename src/tir/{schedule => ir}/block_scope.cc (89%)
create mode 100644 tests/cpp-runtime/opencl/clml_memory_planner.cc
create mode 100644 tests/python/relax/test_frontend_onnx.py
create mode 100644 tests/python/unittest/test_tir_stmt_functor_substitute.py
