You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by tq...@apache.org on 2025/02/15 01:02:43 UTC
(tvm) 01/06: cleanup components related to relay
This is an automated email from the ASF dual-hosted git repository.
tqchen pushed a commit to branch refactor
in repository https://gitbox.apache.org/repos/asf/tvm.git
commit 2c6f4db7b15472448e844eac25d6bfeba17058a4
Author: tqchen <ti...@gmail.com>
AuthorDate: Fri Feb 14 17:27:05 2025 -0500
cleanup components related to relay
---
ci/jenkins/generated/arm_jenkinsfile.groovy | 126 +---
ci/jenkins/generated/gpu_jenkinsfile.groovy | 174 +----
ci/jenkins/templates/arm_jenkinsfile.groovy.j2 | 22 -
ci/jenkins/templates/gpu_jenkinsfile.groovy.j2 | 15 -
jvm/core/src/test/scripts/test_graph_executor.py | 69 --
tests/cpp/relay/backend/aot/aot_lower_main_test.cc | 63 --
tests/cpp/relay/backend/executor_test.cc | 77 ---
tests/cpp/relay/backend/graph_plan_token_alloc.cc | 540 ----------------
tests/cpp/relay/backend/runtime_test.cc | 77 ---
.../cpp/relay/collage/candidate_partition_test.cc | 221 -------
tests/cpp/relay/collage/partition_rule_test.cc | 715 ---------------------
tests/cpp/relay/df_pattern_rewrite_test.cc | 100 ---
tests/cpp/relay/ir/indexed_graph_test.cc | 205 ------
tests/cpp/relay/op/memory/on_device_test.cc | 89 ---
tests/cpp/relay/transforms/device_domains_test.cc | 76 ---
tests/cpp/relay/with_fields_test.cc | 215 -------
tests/cpp/relay_build_module_test.cc | 193 ------
tests/cpp/relay_dismantler_test.cc | 164 -----
tests/cpp/relay_pass_type_infer_test.cc | 43 --
tests/cpp/relay_text_printer_test.cc | 58 --
tests/cpp/relay_transform_sequential_test.cc | 135 ----
.../contrib/test_arm_compute_lib/__init__.py | 17 -
.../contrib/test_arm_compute_lib/infrastructure.py | 341 ----------
.../contrib/test_arm_compute_lib/test_add.py | 168 -----
.../test_arm_compute_lib/test_concatenate.py | 159 -----
.../contrib/test_arm_compute_lib/test_config.json | 8 -
.../contrib/test_arm_compute_lib/test_conv2d.py | 672 -------------------
.../contrib/test_arm_compute_lib/test_dense.py | 430 -------------
.../contrib/test_arm_compute_lib/test_maximum.py | 105 ---
.../contrib/test_arm_compute_lib/test_network.py | 231 -------
.../contrib/test_arm_compute_lib/test_pooling.py | 419 ------------
.../contrib/test_arm_compute_lib/test_reshape.py | 111 ----
.../contrib/test_arm_compute_lib/test_runtime.py | 134 ----
tests/python/runtime/test_runtime_rpc.py | 28 +-
tests/scripts/ci.py | 1 -
tests/scripts/task_java_unittest.sh | 1 -
tests/scripts/task_python_arm_compute_library.sh | 1 -
tests/scripts/task_python_integration.sh | 9 -
tests/scripts/task_python_topi.sh | 8 -
tests/scripts/task_python_unittest.sh | 4 +-
40 files changed, 9 insertions(+), 6215 deletions(-)
diff --git a/ci/jenkins/generated/arm_jenkinsfile.groovy b/ci/jenkins/generated/arm_jenkinsfile.groovy
index 0eab7ac6cb..a8949af697 100644
--- a/ci/jenkins/generated/arm_jenkinsfile.groovy
+++ b/ci/jenkins/generated/arm_jenkinsfile.groovy
@@ -60,7 +60,7 @@
// 'python3 jenkins/generate.py'
// Note: This timestamp is here to ensure that updates to the Jenkinsfile are
// always rebased on main before merging:
-// Generated at 2025-02-09T12:21:01.773823
+// Generated at 2025-02-14T17:21:20.745906
import org.jenkinsci.plugins.pipeline.modeldefinition.Utils
// These are set at runtime from data in ci/jenkins/docker-images.yml, update
@@ -725,108 +725,6 @@ def shard_run_integration_aarch64_4_of_4(node_type) {
-def shard_run_topi_aarch64_1_of_2(node_type) {
- echo 'Begin running on node_type ' + node_type
- if (!skip_ci && is_docs_only_build != 1) {
- node(node_type) {
- ws("workspace/exec_${env.EXECUTOR_NUMBER}/tvm/ut-python-arm") {
- // NOTE: if exception happens, it will be caught outside
- init_git()
- docker_init(ci_arm)
- timeout(time: max_time, unit: 'MINUTES') {
- withEnv([
- 'PLATFORM=arm',
- 'TEST_STEP_NAME=topi: aarch64',
- 'TVM_NUM_SHARDS=2',
- 'TVM_SHARD_INDEX=0',
- "SKIP_SLOW_TESTS=${skip_slow_tests}"], {
- sh(
- script: "./${jenkins_scripts_root}/s3.py --action download --bucket ${s3_bucket} --prefix ${s3_prefix}/arm",
- label: 'Download artifacts from S3',
- )
-
- ci_setup(ci_arm)
- cpp_unittest(ci_arm)
- sh (
- script: "${docker_run} ${ci_arm} ./tests/scripts/task_python_arm_compute_library.sh",
- label: 'Run test_arm_compute_lib test',
- )
- sh (
- script: "${docker_run} ${ci_arm} ./tests/scripts/task_python_topi.sh",
- label: 'Run TOPI tests',
- )
- })
- }
- // only run upload if things are successful
- try {
- sh(
- script: "./${jenkins_scripts_root}/s3.py --action upload --bucket ${s3_bucket} --prefix ${s3_prefix}/pytest-results/topi_aarch64 --items build/pytest-results",
- label: 'Upload JUnits to S3',
- )
-
- junit 'build/pytest-results/*.xml'
- } catch (Exception e) {
- echo 'Exception during JUnit upload: ' + e.toString()
- }
- }
- }
- echo 'End running on node_type ' + node_type
- } else {
- Utils.markStageSkippedForConditional('topi: aarch64 1 of 2')
- }
-}
-
-def shard_run_topi_aarch64_2_of_2(node_type) {
- echo 'Begin running on node_type ' + node_type
- if (!skip_ci && is_docs_only_build != 1) {
- node(node_type) {
- ws("workspace/exec_${env.EXECUTOR_NUMBER}/tvm/ut-python-arm") {
- // NOTE: if exception happens, it will be caught outside
- init_git()
- docker_init(ci_arm)
- timeout(time: max_time, unit: 'MINUTES') {
- withEnv([
- 'PLATFORM=arm',
- 'TEST_STEP_NAME=topi: aarch64',
- 'TVM_NUM_SHARDS=2',
- 'TVM_SHARD_INDEX=1',
- "SKIP_SLOW_TESTS=${skip_slow_tests}"], {
- sh(
- script: "./${jenkins_scripts_root}/s3.py --action download --bucket ${s3_bucket} --prefix ${s3_prefix}/arm",
- label: 'Download artifacts from S3',
- )
-
- ci_setup(ci_arm)
- sh (
- script: "${docker_run} ${ci_arm} ./tests/scripts/task_python_arm_compute_library.sh",
- label: 'Run test_arm_compute_lib test',
- )
- sh (
- script: "${docker_run} ${ci_arm} ./tests/scripts/task_python_topi.sh",
- label: 'Run TOPI tests',
- )
- })
- }
- // only run upload if things are successful
- try {
- sh(
- script: "./${jenkins_scripts_root}/s3.py --action upload --bucket ${s3_bucket} --prefix ${s3_prefix}/pytest-results/topi_aarch64 --items build/pytest-results",
- label: 'Upload JUnits to S3',
- )
-
- junit 'build/pytest-results/*.xml'
- } catch (Exception e) {
- echo 'Exception during JUnit upload: ' + e.toString()
- }
- }
- }
- echo 'End running on node_type ' + node_type
- } else {
- Utils.markStageSkippedForConditional('topi: aarch64 2 of 2')
- }
-}
-
-
def test() {
stage('Test') {
environment {
@@ -877,28 +775,6 @@ def test() {
shard_run_integration_aarch64_4_of_4('ARM-GRAVITON3')
}
},
- 'topi: aarch64 1 of 2': {
- try {
- shard_run_topi_aarch64_1_of_2('ARM-GRAVITON3-SPOT')
- } catch (Throwable ex) {
- // mark the current stage as success
- // and try again via on demand node
- echo 'Exception during SPOT run ' + ex.toString() + ' retry on-demand'
- currentBuild.result = 'SUCCESS'
- shard_run_topi_aarch64_1_of_2('ARM-GRAVITON3')
- }
- },
- 'topi: aarch64 2 of 2': {
- try {
- shard_run_topi_aarch64_2_of_2('ARM-GRAVITON3-SPOT')
- } catch (Throwable ex) {
- // mark the current stage as success
- // and try again via on demand node
- echo 'Exception during SPOT run ' + ex.toString() + ' retry on-demand'
- currentBuild.result = 'SUCCESS'
- shard_run_topi_aarch64_2_of_2('ARM-GRAVITON3')
- }
- },
)
}
}
diff --git a/ci/jenkins/generated/gpu_jenkinsfile.groovy b/ci/jenkins/generated/gpu_jenkinsfile.groovy
index 8c169ffd36..89d7784705 100644
--- a/ci/jenkins/generated/gpu_jenkinsfile.groovy
+++ b/ci/jenkins/generated/gpu_jenkinsfile.groovy
@@ -60,7 +60,7 @@
// 'python3 jenkins/generate.py'
// Note: This timestamp is here to ensure that updates to the Jenkinsfile are
// always rebased on main before merging:
-// Generated at 2025-02-09T12:21:01.812601
+// Generated at 2025-02-14T17:21:20.702568
import org.jenkinsci.plugins.pipeline.modeldefinition.Utils
// These are set at runtime from data in ci/jenkins/docker-images.yml, update
@@ -719,145 +719,6 @@ def shard_run_unittest_GPU_3_of_3(node_type) {
-def shard_run_topi_GPU_1_of_3(node_type) {
- echo 'Begin running on node_type ' + node_type
- if (!skip_ci && is_docs_only_build != 1) {
- node(node_type) {
- ws("workspace/exec_${env.EXECUTOR_NUMBER}/tvm/topi-python-gpu") {
- // NOTE: if exception happens, it will be caught outside
- init_git()
- docker_init(ci_gpu)
- timeout(time: max_time, unit: 'MINUTES') {
- withEnv([
- 'PLATFORM=gpu',
- 'TEST_STEP_NAME=topi: GPU',
- 'TVM_NUM_SHARDS=3',
- 'TVM_SHARD_INDEX=0',
- "SKIP_SLOW_TESTS=${skip_slow_tests}"], {
- sh(
- script: "./${jenkins_scripts_root}/s3.py --action download --bucket ${s3_bucket} --prefix ${s3_prefix}/gpu",
- label: 'Download artifacts from S3',
- )
-
- ci_setup(ci_gpu)
- sh (
- script: "${docker_run} ${ci_gpu} ./tests/scripts/task_python_topi.sh",
- label: 'Run TOPI tests',
- )
- })
- }
- // only run upload if things are successful
- try {
- sh(
- script: "./${jenkins_scripts_root}/s3.py --action upload --bucket ${s3_bucket} --prefix ${s3_prefix}/pytest-results/topi_GPU --items build/pytest-results",
- label: 'Upload JUnits to S3',
- )
-
- junit 'build/pytest-results/*.xml'
- } catch (Exception e) {
- echo 'Exception during JUnit upload: ' + e.toString()
- }
- }
- }
- echo 'End running on node_type ' + node_type
- } else {
- Utils.markStageSkippedForConditional('topi: GPU 1 of 3')
- }
-}
-
-def shard_run_topi_GPU_2_of_3(node_type) {
- echo 'Begin running on node_type ' + node_type
- if (!skip_ci && is_docs_only_build != 1) {
- node(node_type) {
- ws("workspace/exec_${env.EXECUTOR_NUMBER}/tvm/topi-python-gpu") {
- // NOTE: if exception happens, it will be caught outside
- init_git()
- docker_init(ci_gpu)
- timeout(time: max_time, unit: 'MINUTES') {
- withEnv([
- 'PLATFORM=gpu',
- 'TEST_STEP_NAME=topi: GPU',
- 'TVM_NUM_SHARDS=3',
- 'TVM_SHARD_INDEX=1',
- "SKIP_SLOW_TESTS=${skip_slow_tests}"], {
- sh(
- script: "./${jenkins_scripts_root}/s3.py --action download --bucket ${s3_bucket} --prefix ${s3_prefix}/gpu",
- label: 'Download artifacts from S3',
- )
-
- ci_setup(ci_gpu)
- sh (
- script: "${docker_run} ${ci_gpu} ./tests/scripts/task_python_topi.sh",
- label: 'Run TOPI tests',
- )
- })
- }
- // only run upload if things are successful
- try {
- sh(
- script: "./${jenkins_scripts_root}/s3.py --action upload --bucket ${s3_bucket} --prefix ${s3_prefix}/pytest-results/topi_GPU --items build/pytest-results",
- label: 'Upload JUnits to S3',
- )
-
- junit 'build/pytest-results/*.xml'
- } catch (Exception e) {
- echo 'Exception during JUnit upload: ' + e.toString()
- }
- }
- }
- echo 'End running on node_type ' + node_type
- } else {
- Utils.markStageSkippedForConditional('topi: GPU 2 of 3')
- }
-}
-
-def shard_run_topi_GPU_3_of_3(node_type) {
- echo 'Begin running on node_type ' + node_type
- if (!skip_ci && is_docs_only_build != 1) {
- node(node_type) {
- ws("workspace/exec_${env.EXECUTOR_NUMBER}/tvm/topi-python-gpu") {
- // NOTE: if exception happens, it will be caught outside
- init_git()
- docker_init(ci_gpu)
- timeout(time: max_time, unit: 'MINUTES') {
- withEnv([
- 'PLATFORM=gpu',
- 'TEST_STEP_NAME=topi: GPU',
- 'TVM_NUM_SHARDS=3',
- 'TVM_SHARD_INDEX=2',
- "SKIP_SLOW_TESTS=${skip_slow_tests}"], {
- sh(
- script: "./${jenkins_scripts_root}/s3.py --action download --bucket ${s3_bucket} --prefix ${s3_prefix}/gpu",
- label: 'Download artifacts from S3',
- )
-
- ci_setup(ci_gpu)
- sh (
- script: "${docker_run} ${ci_gpu} ./tests/scripts/task_python_topi.sh",
- label: 'Run TOPI tests',
- )
- })
- }
- // only run upload if things are successful
- try {
- sh(
- script: "./${jenkins_scripts_root}/s3.py --action upload --bucket ${s3_bucket} --prefix ${s3_prefix}/pytest-results/topi_GPU --items build/pytest-results",
- label: 'Upload JUnits to S3',
- )
-
- junit 'build/pytest-results/*.xml'
- } catch (Exception e) {
- echo 'Exception during JUnit upload: ' + e.toString()
- }
- }
- }
- echo 'End running on node_type ' + node_type
- } else {
- Utils.markStageSkippedForConditional('topi: GPU 3 of 3')
- }
-}
-
-
def shard_run_docs_GPU_1_of_1(node_type) {
echo 'Begin running on node_type ' + node_type
@@ -955,39 +816,6 @@ def test() {
shard_run_unittest_GPU_3_of_3('GPU')
}
},
- 'topi: GPU 1 of 3': {
- try {
- shard_run_topi_GPU_1_of_3('GPU-SPOT')
- } catch (Throwable ex) {
- // mark the current stage as success
- // and try again via on demand node
- echo 'Exception during SPOT run ' + ex.toString() + ' retry on-demand'
- currentBuild.result = 'SUCCESS'
- shard_run_topi_GPU_1_of_3('GPU')
- }
- },
- 'topi: GPU 2 of 3': {
- try {
- shard_run_topi_GPU_2_of_3('GPU-SPOT')
- } catch (Throwable ex) {
- // mark the current stage as success
- // and try again via on demand node
- echo 'Exception during SPOT run ' + ex.toString() + ' retry on-demand'
- currentBuild.result = 'SUCCESS'
- shard_run_topi_GPU_2_of_3('GPU')
- }
- },
- 'topi: GPU 3 of 3': {
- try {
- shard_run_topi_GPU_3_of_3('GPU-SPOT')
- } catch (Throwable ex) {
- // mark the current stage as success
- // and try again via on demand node
- echo 'Exception during SPOT run ' + ex.toString() + ' retry on-demand'
- currentBuild.result = 'SUCCESS'
- shard_run_topi_GPU_3_of_3('GPU')
- }
- },
'docs: GPU 1 of 1': {
try {
shard_run_docs_GPU_1_of_1('GPU-SPOT')
diff --git a/ci/jenkins/templates/arm_jenkinsfile.groovy.j2 b/ci/jenkins/templates/arm_jenkinsfile.groovy.j2
index f869509105..aa999408a7 100644
--- a/ci/jenkins/templates/arm_jenkinsfile.groovy.j2
+++ b/ci/jenkins/templates/arm_jenkinsfile.groovy.j2
@@ -53,27 +53,5 @@
)
{% endcall %}
-{% call(shard_index, num_shards) m.sharded_test_step(
- name="topi: aarch64",
- ws="tvm/ut-python-arm",
- platform="arm",
- docker_image="ci_arm",
- num_shards=2,
- test_method_names=test_method_names,
-) %}
- {{ m.download_artifacts(tag='arm') }}
- ci_setup(ci_arm)
- {% if shard_index == 1 %}
- cpp_unittest(ci_arm)
- {% endif %}
- sh (
- script: "${docker_run} ${ci_arm} ./tests/scripts/task_python_arm_compute_library.sh",
- label: 'Run test_arm_compute_lib test',
- )
- sh (
- script: "${docker_run} ${ci_arm} ./tests/scripts/task_python_topi.sh",
- label: 'Run TOPI tests',
- )
-{% endcall %}
{{ m.invoke_tests(node="ARM-GRAVITON3", test_method_names=test_method_names) -}}
diff --git a/ci/jenkins/templates/gpu_jenkinsfile.groovy.j2 b/ci/jenkins/templates/gpu_jenkinsfile.groovy.j2
index e57e8c0433..0adb492b98 100644
--- a/ci/jenkins/templates/gpu_jenkinsfile.groovy.j2
+++ b/ci/jenkins/templates/gpu_jenkinsfile.groovy.j2
@@ -87,21 +87,6 @@
)
{% endcall %}
-{% call(shard_index, num_shards) m.sharded_test_step(
- name="topi: GPU",
- num_shards=3,
- ws="tvm/topi-python-gpu",
- platform="gpu",
- docker_image="ci_gpu",
- test_method_names=test_method_names,
-) %}
- {{ m.download_artifacts(tag='gpu') }}
- ci_setup(ci_gpu)
- sh (
- script: "${docker_run} ${ci_gpu} ./tests/scripts/task_python_topi.sh",
- label: 'Run TOPI tests',
- )
-{% endcall %}
{% call(shard_index, num_shards) m.sharded_test_step(
name="docs: GPU",
diff --git a/jvm/core/src/test/scripts/test_graph_executor.py b/jvm/core/src/test/scripts/test_graph_executor.py
deleted file mode 100644
index 676b008205..0000000000
--- a/jvm/core/src/test/scripts/test_graph_executor.py
+++ /dev/null
@@ -1,69 +0,0 @@
-# Licensed to the Apache Software Foundation (ASF) under one
-# or more contributor license agreements. See the NOTICE file
-# distributed with this work for additional information
-# regarding copyright ownership. The ASF licenses this file
-# to you under the Apache License, Version 2.0 (the
-# "License"); you may not use this file except in compliance
-# with the License. You may obtain a copy of the License at
-#
-# http://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing,
-# software distributed under the License is distributed on an
-# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
-# KIND, either express or implied. See the License for the
-# specific language governing permissions and limitations
-# under the License.
-import os
-
-import tvm
-from tvm import te
-import json
-from tvm.contrib import graph_executor
-
-
-def dump_graph_lib(target_dir):
- dim = 4
- A = te.placeholder((dim,), name="A")
- B = te.compute(A.shape, lambda *i: A(*i) + 1.0, name="B")
- sched = te.create_schedule(B.op)
-
- node0 = {"op": "null", "name": "x", "inputs": []}
- node1 = {
- "op": "tvm_op",
- "name": "add",
- "inputs": [[0, 0, 0]],
- "attrs": {"func_name": "myadd", "flatten_data": "1", "num_inputs": "1", "num_outputs": "1"},
- }
- nodes = [node0, node1]
- arg_nodes = [0]
- node_row_ptr = [0, 1, 2]
- outputs = [[1, 0, 0]]
- shape = (4,)
- attrs = {
- "shape": ["list_shape", [shape, shape]],
- "dltype": ["list_str", ["float32", "float32"]],
- "storage_id": ["list_int", [0, 1]],
- }
- graph = {
- "nodes": nodes,
- "arg_nodes": arg_nodes,
- "node_row_ptr": node_row_ptr,
- "heads": outputs,
- "attrs": attrs,
- }
-
- graph = json.dumps(graph)
- mlib = tvm.build(sched, [A, B], "llvm", name="myadd")
-
- mlib.export_library(os.path.join(target_dir, "graph_addone_lib.so"))
- with open(os.path.join(target_dir, "graph_addone.json"), "w") as fo:
- fo.write(graph)
-
-
-if __name__ == "__main__":
- import sys
-
- if len(sys.argv) != 2:
- sys.exit(-1)
- dump_graph_lib(sys.argv[1])
diff --git a/tests/cpp/relay/backend/aot/aot_lower_main_test.cc b/tests/cpp/relay/backend/aot/aot_lower_main_test.cc
deleted file mode 100644
index 0157f031c2..0000000000
--- a/tests/cpp/relay/backend/aot/aot_lower_main_test.cc
+++ /dev/null
@@ -1,63 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include "../../../../../src/relay/backend/aot/aot_lower_main.h"
-
-#include <gtest/gtest.h>
-#include <tvm/relay/parser.h>
-
-namespace tvm {
-namespace relay {
-namespace backend {
-namespace aot {
-
-TEST(AOTLowerMain, ExprAllocatorSkipNestedFunc) {
- constexpr const char* mod_text = R"(
- #[version = "0.0.5"]
- def @main(%x: Tensor[(10, 10), float32]) {
- %0 = fn (%FunctionVar_01: Tensor[(10, 10), float32]) {
- nn.relu(%FunctionVar_01)
- };
- %0(%x)
- }
- )";
- IRModule mod = ParseModule("string", mod_text, {}, {});
- auto host_target = tvm::Target("llvm");
- auto prim_target = tvm::Target(host_target, host_target);
- auto ctxt = tvm::transform::PassContext::Current();
- auto config = tvm::CompilationConfig(ctxt, {prim_target});
- mod = tvm::relay::transform::PlanDevices(config)(mod);
- mod = tvm::relay::transform::InferType()(mod);
-
- StorageMap storage_map;
- std::vector<int> return_sids;
- auto func = Downcast<Function>(mod->Lookup("main"));
- std::tie(storage_map, return_sids) = CreateStorage(func);
-
- auto nested_func = Downcast<Function>(Downcast<Call>(func->body)->op);
- EXPECT_EQ(storage_map.find(nested_func->body), storage_map.end());
- EXPECT_EQ(storage_map.find(nested_func->params[0]), storage_map.end());
- EXPECT_NE(storage_map.find(func->body), storage_map.end());
- EXPECT_NE(storage_map.find(func->params[0]), storage_map.end());
-}
-
-} // namespace aot
-} // namespace backend
-} // namespace relay
-} // namespace tvm
diff --git a/tests/cpp/relay/backend/executor_test.cc b/tests/cpp/relay/backend/executor_test.cc
deleted file mode 100644
index 3367390b27..0000000000
--- a/tests/cpp/relay/backend/executor_test.cc
+++ /dev/null
@@ -1,77 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include <gtest/gtest.h>
-#include <tvm/relay/executor.h>
-
-#include <cmath>
-#include <string>
-namespace tvm {
-namespace relay {
-
-TVM_REGISTER_EXECUTOR("TestExecutor")
- .add_attr_option<Bool>("my_bool")
- .add_attr_option<Array<String>>("your_names")
- .add_attr_option<String>("another_option")
- .add_attr_option<Bool>("defaulty_the_default_option", Bool(false));
-
-TEST(Executor, Create) {
- Map<String, ObjectRef> attrs = {{"my_bool", Bool(true)}};
- Executor my_exec = Executor::Create("TestExecutor", attrs);
- ASSERT_EQ(my_exec->GetAttr<Bool>("my_bool"), true);
- ASSERT_EQ(my_exec->GetAttr<Array<String>>("your_names").defined(), false);
- ASSERT_EQ(my_exec->GetAttr<Bool>("defaulty_the_default_option"), false);
-}
-
-TEST(Executor, UnknownAttr) {
- Map<String, ObjectRef> attrs = {{"woofles", Bool(true)}};
- ASSERT_THROW(Executor::Create("TestExecutor", attrs), Error);
-}
-
-TEST(Executor, IncorrectAttrType) {
- Map<String, ObjectRef> attrs = {{"my_bool", String("snuck_in")}};
- ASSERT_THROW(Executor::Create("TestExecutor", attrs), Error);
-}
-
-TEST(Executor, UnregisteredName) {
- Map<String, ObjectRef> attrs = {};
- ASSERT_THROW(Executor::Create("NeverNameAnExecutorThis", attrs), Error);
-}
-
-TEST(ExecutorRegistry, ListExecutors) {
- Array<String> names = Executor::ListExecutors();
- ICHECK_EQ(names.empty(), false);
- ICHECK_EQ(std::count(std::begin(names), std::end(names), "TestExecutor"), 1);
-}
-
-TEST(ExecutorRegistry, ListExecutorOptions) {
- Map<String, String> attrs = Executor::ListExecutorOptions("TestExecutor");
-
- ICHECK_EQ(attrs.empty(), false);
- ICHECK_EQ(attrs["my_bool"], "IntImm");
- ICHECK_EQ(attrs["your_names"], "Array");
- ICHECK_EQ(attrs["another_option"], "runtime.String");
-}
-
-TEST(ExecutorRegistry, ListExecutorOptionsNoExecutor) {
- ASSERT_THROW(Executor::ListExecutorOptions("NeverNameAnExecutorThis"), Error);
-}
-
-} // namespace relay
-} // namespace tvm
diff --git a/tests/cpp/relay/backend/graph_plan_token_alloc.cc b/tests/cpp/relay/backend/graph_plan_token_alloc.cc
deleted file mode 100644
index 7fca4b26a9..0000000000
--- a/tests/cpp/relay/backend/graph_plan_token_alloc.cc
+++ /dev/null
@@ -1,540 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include <gtest/gtest.h>
-
-#include "../src/relay/backend/token_allocator.h"
-
-namespace tvm {
-namespace relay {
-
-// TokenAllocatorMixed is necessary because in class TokenAllocatorMixed we don't
-// have an access to its protected members. In this class we add new methods
-// which allow us to get and check internal state of class TokenAllocatorMixed
-class TokenAllocatorMixedWrapper : public TokenAllocatorMixed {
- public:
- inline size_t FreeListSize() const { return free_.size(); }
- inline size_t AllocListSize() const { return data_.size(); }
-};
-
-TEST(TokenMixedAlloc, TextureOneToken) {
- TokenAllocatorMixedWrapper alloc;
- int storage_ids = 0;
- EXPECT_EQ(alloc.AllocListSize(), 0);
- EXPECT_EQ(alloc.FreeListSize(), 0);
-
- TensorType tt1({1, 22, 20, 20, 4}, DataType(kDLFloat, 32, 1));
- VirtualDevice vd1(kDLOpenCL, 0, Target("opencl -device=adreno"),
- MemoryScope("global.texture-nhwc"));
- StorageToken tok1 = {
- 1, // ref_counter
- 0, // max bytes
- tt1, // tensor type
- vd1, // virtual device
- -1 // storage_id
- };
- auto size2d = alloc.GetSize2D(&tok1);
- EXPECT_EQ(size2d, 140800);
- EXPECT_EQ(alloc.Request(&tok1), nullptr);
-
- alloc.Alloc(&tok1, storage_ids++);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 0);
-
- tok1.ref_counter -= 1;
- alloc.CheckForRelease(&tok1);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 1);
-}
-
-TEST(TokenMixedAlloc, TextureEqualSizeTokenReuse) {
- TokenAllocatorMixedWrapper alloc;
- int storage_ids = 0;
- EXPECT_EQ(alloc.AllocListSize(), 0);
- EXPECT_EQ(alloc.FreeListSize(), 0);
-
- TensorType tt1({1, 22, 20, 20, 4}, DataType(kDLFloat, 32, 1));
- VirtualDevice vd1(kDLOpenCL, 0, Target("opencl -device=adreno"),
- MemoryScope("global.texture-nhwc"));
- StorageToken tok1 = {
- 1, // ref_counter
- 0, // max bytes
- tt1, // tensor type
- vd1, // virtual device
- -1 // storage_id
- };
- auto size2d = alloc.GetSize2D(&tok1);
- EXPECT_EQ(size2d, 140800);
- EXPECT_EQ(alloc.Request(&tok1), nullptr);
-
- alloc.Alloc(&tok1, storage_ids++);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 0);
-
- tok1.ref_counter -= 1;
- alloc.CheckForRelease(&tok1);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 1);
-
- StorageToken tok2 = {
- 1, // ref_counter
- 0, // max bytes
- tt1, // tensor type
- vd1, // virtual device
- -1 // storage_id
- };
- auto req = alloc.Request(&tok2);
- EXPECT_NE(req, nullptr);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 0);
- EXPECT_EQ(req->storage_id, storage_ids - 1);
- EXPECT_EQ(req->ref_counter, 1);
- auto sizeReq = alloc.GetSize2D(req);
- EXPECT_EQ(sizeReq, 140800);
-
- req->ref_counter -= 1;
- alloc.CheckForRelease(req);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 1);
-
- // Try reuse of the texture memory for buffer object
- VirtualDevice vd2(kDLOpenCL, 0, Target("opencl -device=adreno"), MemoryScope("global"));
- StorageToken tok3 = {
- 1, // ref_counter
- 0, // max bytes
- tt1, // tensor type
- vd2, // virtual device
- -1 // storage_id
- };
- auto req1 = alloc.Request(&tok3);
- EXPECT_NE(req1, nullptr);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 0);
- EXPECT_EQ(req1->storage_id, storage_ids - 1);
- EXPECT_EQ(req1->ref_counter, 1);
- sizeReq = alloc.GetSize2D(req1);
- EXPECT_EQ(sizeReq, 140800);
-
- req1->ref_counter -= 1;
- alloc.CheckForRelease(req1);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 1);
-}
-
-TEST(TokenMixedAlloc, TextureEqualSizeDiffTypes) {
- TokenAllocatorMixedWrapper alloc;
- int storage_ids = 0;
- EXPECT_EQ(alloc.AllocListSize(), 0);
- EXPECT_EQ(alloc.FreeListSize(), 0);
-
- TensorType tt1({1, 22, 20, 20, 4}, DataType(kDLFloat, 32, 1));
- VirtualDevice vd1(kDLOpenCL, 0, Target("opencl -device=adreno"),
- MemoryScope("global.texture-nhwc"));
- StorageToken tok1 = {
- 1, // ref_counter
- 0, // max bytes
- tt1, // tensor type
- vd1, // virtual device
- -1 // storage_id
- };
- auto sizeReq = alloc.GetSize2D(&tok1);
- EXPECT_EQ(sizeReq, 140800);
- EXPECT_EQ(alloc.Request(&tok1), nullptr);
-
- alloc.Alloc(&tok1, storage_ids++);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 0);
-
- tok1.ref_counter -= 1;
- alloc.CheckForRelease(&tok1);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 1);
-
- TensorType tt2({1, 22, 20, 20, 4}, DataType(kDLFloat, 16, 1));
- StorageToken tok2 = {
- 1, // ref_counter
- 0, // max bytes
- tt2, // tensor type
- vd1, // virtual device
- -1 // storage_id
- };
-
- auto req1 = alloc.Request(&tok2);
- EXPECT_NE(req1, nullptr);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 0);
-
- req1->ref_counter -= 1;
- alloc.CheckForRelease(req1);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 1);
-}
-
-TEST(TokenMixedAlloc, TextureDifferentSizesTokenReuse) {
- TokenAllocatorMixedWrapper alloc;
- int storage_ids = 0;
- EXPECT_EQ(alloc.AllocListSize(), 0);
- EXPECT_EQ(alloc.FreeListSize(), 0);
-
- TensorType tt1({1, 22, 20, 20, 4}, DataType(kDLFloat, 32, 1));
- VirtualDevice vd1(kDLOpenCL, 0, Target("opencl -device=adreno"),
- MemoryScope("global.texture-nhwc"));
- StorageToken tok1 = {
- 1, // ref_counter
- 0, // max bytes
- tt1, // tensor type
- vd1, // virtual device
- -1 // storage_id
- };
- auto sizeReq = alloc.GetSize2D(&tok1);
- EXPECT_EQ(sizeReq, 140800);
- EXPECT_EQ(alloc.Request(&tok1), nullptr);
-
- alloc.Alloc(&tok1, storage_ids++);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 0);
-
- tok1.ref_counter -= 1;
- alloc.CheckForRelease(&tok1);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 1);
-
- TensorType tt2({1, 40, 30, 30, 4}, DataType(kDLFloat, 32, 1));
- StorageToken tok2 = {
- 1, // ref_counter
- 0, // max bytes
- tt2, // tensor type
- vd1, // virtual device
- -1 // storage_id
- };
- auto req = alloc.Request(&tok2);
- EXPECT_NE(req, nullptr);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 0);
- EXPECT_EQ(req->storage_id, storage_ids - 1);
- EXPECT_EQ(req->ref_counter, 1);
- sizeReq = alloc.GetSize2D(req);
- EXPECT_EQ(sizeReq, 576000);
-
- req->ref_counter -= 1;
- alloc.CheckForRelease(req);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 1);
-
- TensorType tt3({1, 25, 30, 30, 4}, DataType(kDLFloat, 32, 1));
- StorageToken tok3 = {
- 1, // ref_counter
- 0, // max bytes
- tt3, // tensor type
- vd1, // virtual device
- -1 // storage_id
- };
- auto req2 = alloc.Request(&tok3);
- EXPECT_NE(req2, nullptr);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 0);
- EXPECT_EQ(req2->storage_id, storage_ids - 1);
- EXPECT_EQ(req2->ref_counter, 1);
- sizeReq = alloc.GetSize2D(req2);
- EXPECT_EQ(sizeReq, 576000);
-}
-
-TEST(TokenMixedAlloc, TextureDifferentSizesTokenReuse2) {
- TokenAllocatorMixedWrapper alloc;
- int storage_ids = 0;
- EXPECT_EQ(alloc.AllocListSize(), 0);
- EXPECT_EQ(alloc.FreeListSize(), 0);
-
- TensorType tt1({1, 22, 20, 20, 4}, DataType(kDLFloat, 32, 1));
- VirtualDevice vd1(kDLOpenCL, 0, Target("opencl -device=adreno"),
- MemoryScope("global.texture-nhwc"));
- StorageToken tok1 = {
- 1, // ref_counter
- 0, // max bytes
- tt1, // tensor type
- vd1, // virtual device
- -1 // storage_id
- };
- auto sizeReq = alloc.GetSize2D(&tok1);
- EXPECT_EQ(sizeReq, 140800);
- EXPECT_EQ(alloc.Request(&tok1), nullptr);
-
- alloc.Alloc(&tok1, storage_ids++);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 0);
-
- tok1.ref_counter -= 1;
- alloc.CheckForRelease(&tok1);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 1);
-
- TensorType tt2({1, 5, 30, 20, 4}, DataType(kDLFloat, 32, 1));
- StorageToken tok2 = {
- 1, // ref_counter
- 0, // max bytes
- tt2, // tensor type
- vd1, // virtual device
- -1 // storage_id
- };
- auto req = alloc.Request(&tok2);
- EXPECT_NE(req, nullptr);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 0);
- EXPECT_EQ(req->storage_id, storage_ids - 1);
- EXPECT_EQ(req->ref_counter, 1);
- sizeReq = alloc.GetSize2D(req);
- EXPECT_EQ(sizeReq, 140800);
-}
-
-TEST(TokenMixedAlloc, TextureSameSizesButDiffMemoryScopes) {
- TokenAllocatorMixedWrapper alloc;
- int storage_ids = 0;
- EXPECT_EQ(alloc.AllocListSize(), 0);
- EXPECT_EQ(alloc.FreeListSize(), 0);
-
- TensorType tt1({28, 676, 1, 1, 4}, DataType(kDLFloat, 32, 1));
- VirtualDevice vd1(kDLOpenCL, 0, Target("opencl -device=adreno"),
- MemoryScope("global.texture-weight"));
- StorageToken tok1 = {
- 1, // ref_counter
- 0, // max bytes
- tt1, // tensor type
- vd1, // virtual device
- -1 // storage_id
- };
- auto sizeReq = alloc.GetSize2D(&tok1);
- EXPECT_EQ(sizeReq, 302848);
- EXPECT_EQ(alloc.Request(&tok1), nullptr);
-
- alloc.Alloc(&tok1, storage_ids++);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 0);
-
- tok1.ref_counter -= 1;
- alloc.CheckForRelease(&tok1);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 1);
-
- TensorType tt2({1, 28, 26, 26, 4}, DataType(kDLFloat, 32, 1));
- VirtualDevice vd2(kDLOpenCL, 0, Target("opencl -device=adreno"),
- MemoryScope("global.texture-nhwc"));
- StorageToken tok2 = {
- 1, // ref_counter
- 0, // max bytes
- tt2, // tensor type
- vd2, // virtual device
- -1 // storage_id
- };
- auto tok2Size = alloc.GetSize2D(&tok2);
- EXPECT_EQ(tok2Size, 302848);
-
- auto req = alloc.Request(&tok2);
- EXPECT_NE(req, nullptr);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 0);
-
- req->ref_counter -= 1;
- alloc.CheckForRelease(req);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 1);
-}
-
-TEST(TokenMixedAlloc, OneToken) {
- TokenAllocatorMixedWrapper alloc;
- int storage_ids = 0;
- EXPECT_EQ(alloc.AllocListSize(), 0);
- EXPECT_EQ(alloc.FreeListSize(), 0);
-
- TensorType tt1({1, 22, 20, 20, 4}, DataType(kDLFloat, 32, 1));
- VirtualDevice vd1(kDLOpenCL, 0, Target("opencl"));
- StorageToken tok1 = {
- 1, // ref_counter
- 0, // max bytes
- tt1, // tensor type
- vd1, // virtual device
- -1 // storage_id
- };
- EXPECT_EQ(alloc.Request(&tok1), nullptr);
-
- alloc.Alloc(&tok1, storage_ids++);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 0);
-
- tok1.ref_counter -= 1;
- alloc.CheckForRelease(&tok1);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 1);
-}
-
-TEST(TokenMixedAlloc, EqualSizeTokenReuse) {
- TokenAllocatorMixedWrapper alloc;
- int storage_ids = 0;
- EXPECT_EQ(alloc.AllocListSize(), 0);
- EXPECT_EQ(alloc.FreeListSize(), 0);
-
- TensorType tt1({1, 22, 20, 20, 4}, DataType(kDLFloat, 32, 1));
- VirtualDevice vd1(kDLOpenCL, 0, Target("opencl"));
- StorageToken tok1 = {
- 1, // ref_counter
- 0, // max bytes
- tt1, // tensor type
- vd1, // virtual device
- -1 // storage_id
- };
- EXPECT_EQ(alloc.Request(&tok1), nullptr);
-
- alloc.Alloc(&tok1, storage_ids++);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 0);
-
- tok1.ref_counter -= 1;
- alloc.CheckForRelease(&tok1);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 1);
-
- StorageToken tok2 = {
- 1, // ref_counter
- 0, // max bytes
- tt1, // tensor type
- vd1, // virtual device
- -1 // storage_id
- };
- auto req = alloc.Request(&tok2);
- EXPECT_NE(req, nullptr);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 0);
- EXPECT_EQ(req->storage_id, storage_ids - 1);
- EXPECT_EQ(req->ref_counter, 1);
-
- req->ref_counter -= 1;
- alloc.CheckForRelease(req);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 1);
-}
-
-TEST(TokenMixedAlloc, EqualSizeDiffTypes) {
- TokenAllocatorMixedWrapper alloc;
- int storage_ids = 0;
- EXPECT_EQ(alloc.AllocListSize(), 0);
- EXPECT_EQ(alloc.FreeListSize(), 0);
-
- TensorType tt1({1, 22, 20, 20, 4}, DataType(kDLFloat, 32, 1));
- VirtualDevice vd1(kDLOpenCL, 0, Target("opencl"));
- StorageToken tok1 = {
- 1, // ref_counter
- 0, // max bytes
- tt1, // tensor type
- vd1, // virtual device
- -1 // storage_id
- };
- EXPECT_EQ(alloc.Request(&tok1), nullptr);
-
- alloc.Alloc(&tok1, storage_ids++);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 0);
-
- tok1.ref_counter -= 1;
- alloc.CheckForRelease(&tok1);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 1);
-
- TensorType tt2({1, 22, 20, 20, 4}, DataType(kDLFloat, 16, 1));
- StorageToken tok2 = {
- 1, // ref_counter
- 0, // max bytes
- tt2, // tensor type
- vd1, // virtual device
- -1 // storage_id
- };
-
- auto req1 = alloc.Request(&tok2);
- EXPECT_NE(req1, nullptr);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 0);
-
- req1->ref_counter -= 1;
- alloc.CheckForRelease(req1);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 1);
-}
-
-TEST(TokenMixedAlloc, DifferentSizesTokenReuse) {
- TokenAllocatorMixedWrapper alloc;
- int storage_ids = 0;
- EXPECT_EQ(alloc.AllocListSize(), 0);
- EXPECT_EQ(alloc.FreeListSize(), 0);
-
- TensorType tt1({1, 22, 20, 20, 4}, DataType(kDLFloat, 32, 1));
- VirtualDevice vd1(kDLOpenCL, 0, Target("opencl"));
- StorageToken tok1 = {
- 1, // ref_counter
- 0, // max bytes
- tt1, // tensor type
- vd1, // virtual device
- -1 // storage_id
- };
- EXPECT_EQ(alloc.Request(&tok1), nullptr);
-
- alloc.Alloc(&tok1, storage_ids++);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 0);
-
- tok1.ref_counter -= 1;
- alloc.CheckForRelease(&tok1);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 1);
-
- TensorType tt2({1, 40, 30, 30, 4}, DataType(kDLFloat, 32, 1));
- StorageToken tok2 = {
- 1, // ref_counter
- 0, // max bytes
- tt2, // tensor type
- vd1, // virtual device
- -1 // storage_id
- };
- auto req = alloc.Request(&tok2);
- EXPECT_NE(req, nullptr);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 0);
- EXPECT_EQ(req->storage_id, storage_ids - 1);
- EXPECT_EQ(req->ref_counter, 1);
-
- req->ref_counter -= 1;
- alloc.CheckForRelease(req);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 1);
-
- TensorType tt3({1, 25, 30, 30, 4}, DataType(kDLFloat, 32, 1));
- StorageToken tok3 = {
- 1, // ref_counter
- 0, // max bytes
- tt3, // tensor type
- vd1, // virtual device
- -1 // storage_id
- };
- auto req2 = alloc.Request(&tok3);
- EXPECT_NE(req2, nullptr);
- EXPECT_EQ(alloc.AllocListSize(), 1);
- EXPECT_EQ(alloc.FreeListSize(), 0);
- EXPECT_EQ(req2->storage_id, storage_ids - 1);
- EXPECT_EQ(req2->ref_counter, 1);
-}
-
-} // namespace relay
-} // namespace tvm
diff --git a/tests/cpp/relay/backend/runtime_test.cc b/tests/cpp/relay/backend/runtime_test.cc
deleted file mode 100644
index adabb9b9b6..0000000000
--- a/tests/cpp/relay/backend/runtime_test.cc
+++ /dev/null
@@ -1,77 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include <gtest/gtest.h>
-#include <tvm/relay/runtime.h>
-
-#include <cmath>
-#include <string>
-namespace tvm {
-namespace relay {
-
-TVM_REGISTER_RUNTIME("TestRuntime")
- .add_attr_option<runtime::Bool>("my_bool")
- .add_attr_option<Array<String>>("your_names")
- .add_attr_option<String>("another_option")
- .add_attr_option<runtime::Bool>("defaulty_the_default_option", runtime::Bool(false));
-
-TEST(Runtime, Create) {
- Map<String, ObjectRef> attrs = {{"my_bool", runtime::Bool(true)}};
- Runtime my_runtime = Runtime::Create("TestRuntime", attrs);
- ASSERT_EQ(my_runtime->GetAttr<Bool>("my_bool"), true);
- ASSERT_EQ(my_runtime->GetAttr<Array<String>>("your_names").defined(), false);
- ASSERT_EQ(my_runtime->GetAttr<Bool>("defaulty_the_default_option"), false);
-}
-
-TEST(Runtime, UnknownAttr) {
- Map<String, ObjectRef> attrs = {{"woofles", runtime::Bool(true)}};
- ASSERT_THROW(Runtime::Create("TestRuntime", attrs), Error);
-}
-
-TEST(Runtime, IncorrectAttrType) {
- Map<String, ObjectRef> attrs = {{"my_bool", String("snuck_in")}};
- ASSERT_THROW(Runtime::Create("TestRuntime", attrs), Error);
-}
-
-TEST(Runtime, UnregisteredName) {
- Map<String, ObjectRef> attrs = {};
- ASSERT_THROW(Runtime::Create("NeverNameAnRuntimeThis", attrs), Error);
-}
-
-TEST(RuntimeRegistry, ListRuntimes) {
- Array<String> names = Runtime::ListRuntimes();
- ICHECK_EQ(names.empty(), false);
- ICHECK_EQ(std::count(std::begin(names), std::end(names), "TestRuntime"), 1);
-}
-
-TEST(RuntimeRegistry, ListRuntimeOptions) {
- Map<String, String> attrs = Runtime::ListRuntimeOptions("TestRuntime");
-
- ICHECK_EQ(attrs.empty(), false);
- ICHECK_EQ(attrs["my_bool"], "runtime.BoxBool");
- ICHECK_EQ(attrs["your_names"], "Array");
- ICHECK_EQ(attrs["another_option"], "runtime.String");
-}
-
-TEST(RuntimeRegistry, ListRuntimeOptionsNoRuntime) {
- ASSERT_THROW(Runtime::ListRuntimeOptions("NeverNameAnRuntimeThis"), Error);
-}
-
-} // namespace relay
-} // namespace tvm
diff --git a/tests/cpp/relay/collage/candidate_partition_test.cc b/tests/cpp/relay/collage/candidate_partition_test.cc
deleted file mode 100644
index d298a493c1..0000000000
--- a/tests/cpp/relay/collage/candidate_partition_test.cc
+++ /dev/null
@@ -1,221 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include "../../../src/relay/collage/candidate_partition.h"
-
-#include <gtest/gtest.h>
-#include <tvm/relay/expr.h>
-#include <tvm/relay/function.h>
-#include <tvm/relay/parser.h>
-#include <tvm/relay/transform.h>
-
-#include "../../../../src/relay/collage/mock_cost_estimator.h"
-#include "../../../src/relay/collage/partition_spec.h"
-
-namespace tvm {
-namespace relay {
-namespace collage {
-namespace {
-
-// NOTE: CandidatePartition::ParallelRewrite is effectively tested in partition_rule_test.cc
-// so not re-tested here. The only other non-trivial code is CandidatePartition::EstimateCost
-
-Function MakeTestFunction(const std::string& mod_text) {
- IRModule mod = ParseModule("string", mod_text, {}, {});
- mod = transform::CapturePostDfsIndexInSpans()(mod);
- auto func = Downcast<Function>(mod->Lookup("main"));
- LOG(INFO) << "------- input function -------";
- LOG(INFO) << PrettyPrint(func);
- LOG(INFO) << "------------------------------";
- return func;
-}
-
-PartitionSpec StandardSpec() { return PartitionSpec("test_spec", Target("llvm"), {}); }
-
-String AlwaysInvalid(const Function& function) { return "invalid"; }
-
-PartitionSpec AlwaysInvalidSpec() {
- return PartitionSpec("test_spec", Target("llvm"), {}, AlwaysInvalid);
-}
-
-/*!
- * \brief Returns candidate containing nodes with given \p indexes wrapped within a
- * "Primitive" and "Compiler" function.
- */
-CandidatePartition MakeCandidate(const DataflowGraph& graph, const PartitionSpec& spec,
- const std::vector<PostDfsIndex>& indexes) {
- IndexSet inside(graph.size(), indexes);
- SubGraph inner_sub_graph(graph, inside);
- FunctionAttrsMap attrs_map;
- attrs_map.Set(attr::kPrimitive, Integer(1));
- attrs_map.Set(attr::kCompiler, String("llvm"));
- NestedSubGraph nested_sub_graph(inner_sub_graph, attrs_map);
- SubGraph outer_sub_graph(graph, inside, inner_sub_graph->kind_, inner_sub_graph->label_,
- {nested_sub_graph});
- return CandidatePartition(/*rule_name=*/"", outer_sub_graph, spec);
-}
-
-CostEstimator StandardEstimator() {
- Map<String, Integer> target_costs;
- target_costs.Set("llvm", 3);
- return MockCostEstimator(std::move(target_costs));
-}
-
-CostEstimator AlternateEstimator() {
- Map<String, Integer> target_costs;
- target_costs.Set("llvm", 7);
- return MockCostEstimator(std::move(target_costs));
-}
-
-std::shared_ptr<CandidateFunctionCache> Cache() {
- return std::make_shared<CandidateFunctionCache>(std::make_shared<NameSupply>("test"));
-}
-
-TEST(CandidatePartition, EstimateCost_Simple) {
- constexpr const char* kMod = R"(
- #[version = "0.0.5"]
- def @main(%x: Tensor[(10, 10), float32]) {
- %0 = abs(%x); // 3
- %1 = nn.relu(%0); // 4
- nn.relu(%1) // 5
- }
- )";
- auto func = MakeTestFunction(kMod);
- auto graph = DataflowGraph(func);
- auto spec = StandardSpec();
- auto candidate = MakeCandidate(graph, spec, {3, 4});
- auto estimator = StandardEstimator();
- auto cache = Cache();
-
- {
- auto cost = candidate->EstimatedCost(graph, estimator, cache);
- ASSERT_TRUE(cost.is_value());
- // cost is 3 for nn.rulu plus 3 * 0.9 for the nested abs
- ASSERT_EQ(cost.value(), 5.7);
- }
-}
-
-TEST(CandidatePartition, EstimateCost_AlreadyCached) {
- constexpr const char* kMod = R"(
- #[version = "0.0.5"]
- def @main(%x: Tensor[(10, 10), float32]) {
- %0 = abs(%x); // 3
- %1 = nn.relu(%0); // 4
- nn.relu(%1) // 5
- }
- )";
- auto func = MakeTestFunction(kMod);
- auto graph = DataflowGraph(func);
- auto spec = StandardSpec();
- auto candidate = MakeCandidate(graph, spec, {3, 4});
- candidate->cost_ = Cost::Value(42.0);
- auto estimator = StandardEstimator();
- auto cache = Cache();
-
- {
- auto cost = candidate->EstimatedCost(graph, estimator, cache);
- ASSERT_TRUE(cost.is_value());
- ASSERT_EQ(cost.value(), 42.0);
- }
-}
-
-TEST(CandidatePartition, EstimateCost_Invalid) {
- constexpr const char* kMod = R"(
- #[version = "0.0.5"]
- def @main(%x: Tensor[(10, 10), float32]) {
- %0 = abs(%x); // 3
- %1 = nn.relu(%0); // 4
- nn.relu(%1) // 5
- }
- )";
- auto func = MakeTestFunction(kMod);
- auto graph = DataflowGraph(func);
- auto spec = AlwaysInvalidSpec();
- auto candidate = MakeCandidate(graph, spec, {3, 4});
- auto estimator = StandardEstimator();
- auto cache = Cache();
-
- {
- auto cost = candidate->EstimatedCost(graph, estimator, cache);
- ASSERT_TRUE(cost.is_invalid());
- }
-}
-
-TEST(CandidatePartition, EstimateCost_Cached) {
- constexpr const char* kMod = R"(
- #[version = "0.0.5"]
- def @main(%x: Tensor[(10, 10), float32]) {
- %0 = abs(%x); // 4
- %1 = nn.relu(%0); // 5
- %2 = abs(%1); // 6
- %3 = nn.relu(%2); // 7
- add(%1, %3) // 8
- }
- )";
- auto func = MakeTestFunction(kMod);
- auto graph = DataflowGraph(func);
- auto spec = StandardSpec();
- auto candidateA = MakeCandidate(graph, spec, {4, 5});
- auto candidateB = MakeCandidate(graph, spec, {6, 7});
- auto standard_estimator = StandardEstimator();
- auto alternate_estimator = AlternateEstimator();
- auto cache = Cache();
-
- {
- // First candidate estimated as per usual.
- auto costA = candidateA->EstimatedCost(graph, standard_estimator, cache);
- ASSERT_TRUE(costA.is_value());
- ASSERT_EQ(costA.value(), 5.7);
-
- // Second candidate is structurally equal to first, so reuse first's cost even though
- // estimator has different weights.
- auto costB = candidateB->EstimatedCost(graph, alternate_estimator, cache);
- ASSERT_TRUE(costB.is_value());
- ASSERT_EQ(costB.value(), costA.value());
- }
-}
-
-TEST(CandidatePartition, EstimateCost_EtaExpandTuples) {
- constexpr const char* kMod = R"(
- #[version = "0.0.5"]
- def @main(%x: Tensor[(10, 10), float32]) {
- %0 = abs(%x); // 3
- %1 = nn.relu(%0); // 5
- %2 = (%0, %1); // 6
- concatenate(%2) // 7
- }
- )";
- auto func = MakeTestFunction(kMod);
- auto graph = DataflowGraph(func);
- auto spec = StandardSpec();
- auto candidate = MakeCandidate(graph, spec, {7});
- auto estimator = StandardEstimator();
- auto cache = Cache();
-
- {
- auto cost = candidate->EstimatedCost(graph, estimator, cache);
- ASSERT_TRUE(cost.is_value());
- ASSERT_EQ(cost.value(), 3);
- }
-}
-
-} // namespace
-} // namespace collage
-} // namespace relay
-} // namespace tvm
diff --git a/tests/cpp/relay/collage/partition_rule_test.cc b/tests/cpp/relay/collage/partition_rule_test.cc
deleted file mode 100644
index effe0b1fa0..0000000000
--- a/tests/cpp/relay/collage/partition_rule_test.cc
+++ /dev/null
@@ -1,715 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include "../../../src/relay/collage/partition_rule.h"
-
-#include <gtest/gtest.h>
-#include <tvm/relay/expr.h>
-#include <tvm/relay/function.h>
-#include <tvm/relay/parser.h>
-#include <tvm/relay/transform.h>
-
-#include "../../../src/relay/collage/partition_spec.h"
-
-namespace tvm {
-namespace relay {
-namespace collage {
-namespace {
-
-Constant MakeConstant(std::initializer_list<ShapeTuple::index_type> shape) {
- return Constant(runtime::NDArray::Empty(shape, DataType::Float(32), {kDLCPU, 0}));
-}
-
-Function MakeTestFunction(
- const std::string& mod_text,
- const std::initializer_list<std::initializer_list<ShapeTuple::index_type>>& constant_shapes =
- {}) {
- Array<ObjectRef> constants;
- for (const auto& shape : constant_shapes) {
- constants.push_back(MakeConstant(shape));
- }
- Map<String, Array<ObjectRef>> metatable;
- metatable.Set("relay.Constant", constants);
- IRModule mod = ParseModule("string", mod_text, {}, metatable);
- mod = transform::CapturePostDfsIndexInSpans()(mod);
- auto func = Downcast<Function>(mod->Lookup("main"));
- LOG(INFO) << "------- input function -------";
- LOG(INFO) << PrettyPrint(func);
- LOG(INFO) << "------------------------------";
- return func;
-}
-
-Function StandardTestFunction() {
- constexpr const char* kMod = R"(
- #[version = "0.0.5"]
- def @main(%x: Tensor[(10, 10), float32]) {
- // index, kind
- %0 = abs(%x); // 3, E
- %1 = nn.relu(%0); // 4, E
- nn.relu(%1) // 5, E
- }
- )";
- return MakeTestFunction(kMod);
-}
-
-Function VariantTestFunction() {
- constexpr const char* kMod = R"(
- #[version = "0.0.5"]
- def @main(%x: Tensor[(10, 10), float32]) {
- // index, kind
- %0 = abs(%x); // 4, E
- %1 = add(%0, %x); // 5, E
- shape_of(%1) // 6, O
- }
- )";
- return MakeTestFunction(kMod);
-}
-
-Function GPT2ExtractOps() {
- constexpr const char* kMod = R"(
- #[version = "0.0.5"]
- def @main(%x: Tensor[(1600, 768), float32]) {
- // index, kind
- %60 = nn.dense(%x, meta[relay.Constant][0] /*(3072, 768)*/, units=3072); // 6, A
- %61 = add(%60, meta[relay.Constant][1] /*(3072)*/); // 8, B
- %62 = reshape(%61, newshape=[50, 32, 3072]); // 9, I
- %63 = power(%62, 3f); // 15, B
- %64 = multiply(%63, 0.044715f); // 17, B
- %65 = add(%62, %64); // 18, B
- %66 = multiply(%65, 0.797885f); // 20, B
- %67 = tanh(%66); // 21, E
- %68 = multiply(%62, 0.5f); // 11, B
- %69 = add(%67, 1f); // 23, B
- multiply(%68, %69) // 24, B
- }
- )";
- return MakeTestFunction(kMod, {{3072, 768}, {3072}});
-}
-
-Function GPT2ExtractTuples() {
- constexpr const char* kMod = R"(
- #[version = "0.0.5"]
- def @main(%x: Tensor[(50, 32, 2304), float32]) {
- // index, kind
- %19 = split(%x, indices_or_sections=[768, 1536], axis=2); // 6, I
- %23 = %19.1; // 7
- %24 = reshape(%23, newshape=[50, 32, 12, 64]); // 8, I
- %35 = %19.2; // 11
- %36 = reshape(%35, newshape=[50, 32, 12, 64]); // 12, I
- %37 = transpose(%36, axes=[0, 2, 1, 3]); // 13, I
- %855 = transpose(%24, axes=[0, 2, 1, 3]); // 9, I
- %856 = expand_dims(%855, axis=0); // 10, B
- %857 = expand_dims(%37, axis=0); // 14, B
- %858 = (%856, %857); // 15, B
- concatenate(%858) // 16, I
- }
- )";
- return MakeTestFunction(kMod);
-}
-
-PartitionSpec StandardSpec(const std::string& spec_name = "test_spec",
- const std::string& target = "llvm") {
- return PartitionSpec(spec_name, Target(target), {});
-}
-
-std::vector<CandidatePartition> ActualCandidates(const DataflowGraph& graph, const Function& func,
- const PartitionSpec& spec,
- const PartitionRule& rule) {
- auto candidates = rule->AllCandidates(graph, spec);
- LOG(INFO) << "--------- actual candidates -------------";
- for (const auto& candidate : candidates) {
- LOG(INFO) << candidate->ToString();
- }
- LOG(INFO) << "-----------------------------------------";
- return candidates;
-}
-
-std::vector<CandidatePartition> ExpectedCandidates(
- const DataflowGraph& graph, const PartitionSpec& spec,
- const std::vector<std::vector<PostDfsIndex>>& index_sets) {
- std::vector<CandidatePartition> candidate_partitions;
- for (const auto& indexes : index_sets) {
- auto subgraph = SubGraph(graph, IndexSet(graph.size(), indexes));
- auto candidate = CandidatePartition(/*rule_name=*/"", subgraph, spec);
- candidate_partitions.emplace_back(std::move(candidate));
- }
- return candidate_partitions;
-}
-
-void AssertEqual(const std::vector<CandidatePartition>& actual,
- const std::vector<CandidatePartition>& expected) {
- ASSERT_EQ(actual.size(), expected.size());
- std::set<CandidatePartition, CandidatePartitionCompare> actual_set(actual.begin(), actual.end());
- std::set<CandidatePartition, CandidatePartitionCompare> expected_set(expected.begin(),
- expected.end());
- ASSERT_EQ(actual_set.size(), expected_set.size());
- for (const auto& actual_candidate : actual_set) {
- ASSERT_EQ(expected_set.count(actual_candidate), 1) << actual_candidate->ToString();
- }
-}
-
-void AssertEqual(const Expr& actual, const Expr& expected) {
- ASSERT_TRUE(StructuralEqual()(actual, expected)) << PrettyPrint(actual);
-}
-
-TEST(PartitionRule, DFPatternSingleOp) {
- auto func = StandardTestFunction();
- auto graph = DataflowGraph(func);
- auto spec = StandardSpec();
-
- {
- auto pattern = IsOp("nn.relu")({IsWildcard()});
- auto rule = DFPatternPartitionRule("relu_pattern", pattern);
-
- auto actual_candidates = ActualCandidates(graph, func, spec, rule);
-
- auto expected_candidates = ExpectedCandidates(graph, spec, {{4}, {5}});
- AssertEqual(actual_candidates, expected_candidates);
- }
-}
-
-TEST(PartitionRule, DFPatternOverlap) {
- auto func = StandardTestFunction();
- auto graph = DataflowGraph(func);
- auto spec = StandardSpec();
-
- {
- auto pattern =
- IsOp("nn.relu")({IsOp("nn.relu")({IsWildcard()}) || IsOp("abs")({IsWildcard()})});
- auto rule = DFPatternPartitionRule("relu+abs_pattern", pattern);
-
- auto actual_candidates = ActualCandidates(graph, func, spec, rule);
-
- auto expected_candidates = ExpectedCandidates(graph, spec, {{3, 4}, {4, 5}});
- AssertEqual(actual_candidates, expected_candidates);
- }
-}
-
-TEST(PartitionRule, Composite) {
- auto func = StandardTestFunction();
- auto graph = DataflowGraph(func);
- auto spec = StandardSpec();
-
- constexpr const char* kExpectedMod = R"(
- #[version = "0.0.5"]
- def @main(%x: Tensor[(10, 10), float32]) {
- %0 = abs(%x);
- %1 = fn (%FunctionVar_01: Tensor[(10, 10), float32], Composite="composite") {
- nn.relu(%FunctionVar_01)
- };
- %2 = %1(%0);
- %3 = fn (%FunctionVar_0: Tensor[(10, 10), float32], Composite="composite") {
- nn.relu(%FunctionVar_0)
- };
- %3(%2)
- }
- )";
- Expr expected_expr = MakeTestFunction(kExpectedMod);
-
- {
- auto pattern = IsOp("nn.relu")({IsWildcard()});
- auto df_rule = DFPatternPartitionRule("relu_pattern", pattern);
- auto composite_rule = CompositePartitionRule("composite", df_rule);
-
- auto actual_candidates = ActualCandidates(graph, func, spec, composite_rule);
- auto actual_expr = CandidatePartition::ParallelRewrite(graph, actual_candidates);
-
- auto expected_candidates = ExpectedCandidates(graph, spec, {{4}, {5}});
- AssertEqual(actual_candidates, expected_candidates);
- AssertEqual(actual_expr, expected_expr);
- }
-}
-
-TEST(PartitionRule, PrimitiveTVM) {
- auto func = StandardTestFunction();
- auto graph = DataflowGraph(func);
- auto spec = StandardSpec();
-
- constexpr const char* kExpectedMod = R"(
- #[version = "0.0.5"]
- def @main(%x: Tensor[(10, 10), float32]) {
- %0 = abs(%x);
- %1 = fn (%FunctionVar_01: Tensor[(10, 10), float32], Primitive=1) {
- nn.relu(%FunctionVar_01)
- };
- %2 = %1(%0);
- %3 = fn (%FunctionVar_0: Tensor[(10, 10), float32], Primitive=1) {
- nn.relu(%FunctionVar_0)
- };
- %3(%2)
- }
- )";
- Expr expected_expr = MakeTestFunction(kExpectedMod);
-
- {
- auto pattern = IsOp("nn.relu")({IsWildcard()});
- auto df_rule = DFPatternPartitionRule("relu_pattern", pattern);
- auto primitive_rule = PrimitivePartitionRule("primitive", df_rule);
-
- auto actual_candidates = ActualCandidates(graph, func, spec, primitive_rule);
- auto actual_expr = CandidatePartition::ParallelRewrite(graph, actual_candidates);
-
- auto expected_candidates = ExpectedCandidates(graph, spec, {{4}, {5}});
- AssertEqual(actual_candidates, expected_candidates);
- AssertEqual(actual_expr, expected_expr);
- }
-}
-
-TVM_REGISTER_TARGET_KIND("test_ext_codegen", kDLCUDA)
- .set_attr<Bool>(tvm::attr::kIsExternalCodegen, Bool(true));
-
-TEST(PartitionRule, PrimitiveExternal) {
- auto func = StandardTestFunction();
- auto graph = DataflowGraph(func);
- auto spec = StandardSpec("test_ext_codegen", "test_ext_codegen");
-
- constexpr const char* kExpectedMod = R"(
- #[version = "0.0.5"]
- def @main(%x: Tensor[(10, 10), float32]) {
- %0 = abs(%x);
- %1 = fn (%FunctionVar_01: Tensor[(10, 10), float32], Primitive=1, Compiler="test_ext_codegen") {
- nn.relu(%FunctionVar_01)
- };
- %2 = %1(%0);
- %3 = fn (%FunctionVar_0: Tensor[(10, 10), float32], Primitive=1, Compiler="test_ext_codegen") {
- nn.relu(%FunctionVar_0)
- };
- %3(%2)
- }
- )";
- Expr expected_expr = MakeTestFunction(kExpectedMod);
-
- {
- auto pattern = IsOp("nn.relu")({IsWildcard()});
- auto df_rule = DFPatternPartitionRule("relu_pattern", pattern);
- auto primitive_rule = PrimitivePartitionRule("primitive", df_rule);
-
- auto actual_candidates = ActualCandidates(graph, func, spec, primitive_rule);
- auto actual_expr = CandidatePartition::ParallelRewrite(graph, actual_candidates);
-
- auto expected_candidates = ExpectedCandidates(graph, spec, {{4}, {5}});
- AssertEqual(actual_candidates, expected_candidates);
- AssertEqual(actual_expr, expected_expr);
- }
-}
-
-TEST(PartitionRule, Union) {
- auto func = StandardTestFunction();
- auto graph = DataflowGraph(func);
- auto spec = StandardSpec();
-
- {
- auto abs_pattern = IsOp("abs")({IsWildcard()});
- auto abs_rule = DFPatternPartitionRule("abs_pattern", abs_pattern);
- auto relu_pattern = IsOp("nn.relu")({IsWildcard()});
- auto relu_rule = DFPatternPartitionRule("relu_pattern", relu_pattern);
- auto union_rule = UnionPartitionRule("union", {abs_rule, relu_rule});
-
- auto actual_candidates = ActualCandidates(graph, func, spec, union_rule);
-
- auto expected_candidates = ExpectedCandidates(graph, spec, {{3}, {4}, {5}});
- AssertEqual(actual_candidates, expected_candidates);
- }
-}
-
-TEST(PartitionRule, OpCallByKind) {
- auto func = VariantTestFunction();
- auto graph = DataflowGraph(func);
- auto spec = StandardSpec();
-
- {
- auto rule = OpCallByKindPartitionRule("op_call_by_kind");
- auto actual_candidates = ActualCandidates(graph, func, spec, rule);
-
- auto expected_candidates = ExpectedCandidates(graph, spec, {{4}, {5}});
- AssertEqual(actual_candidates, expected_candidates);
- }
-}
-
-TEST(PartitionRule, Combine_ByKind) {
- auto func = GPT2ExtractOps();
- auto graph = DataflowGraph(func);
- auto spec = StandardSpec();
-
- {
- // Prime the system by picking out all 11 calls to non-opaque ops.
- auto sub_rule = OpCallByKindPartitionRule("op_call_by_kind");
- // Combine all <= kOutEWiseFusable (A) actual_candidates (ie anything) with downstream
- // <= kBroadcast (B) actual_candidates (ie B or E).
- Array<SimpleCombinerRule> simple_rules;
- simple_rules.push_back(ByKindSimpleCombinerRule(/*upstream_kind=*/kOutEWiseFusable,
- /*downstream_kind=*/kBroadcast));
- Array<CombinerRule> combiner_rules;
- combiner_rules.push_back(AllSimpleCombinerRule("all_simple", std::move(simple_rules)));
- // Build the overall partition rule.
- auto rule = CombinePartitionRule("combine_by_kind_A_B", std::move(sub_rule),
- std::move(combiner_rules), /*max_depth=*/3);
-
- auto actual_candidates = ActualCandidates(graph, func, spec, rule);
-
- // The original calls.
- std::vector<std::vector<PostDfsIndex>> expected;
- expected.push_back({6});
- expected.push_back({8});
- expected.push_back({9});
- expected.push_back({11});
- expected.push_back({15});
- expected.push_back({17});
- expected.push_back({18});
- expected.push_back({20});
- expected.push_back({21});
- expected.push_back({23});
- expected.push_back({24});
-
- // nn.dense (A) and the following add (B)
- expected.push_back({6, 8});
-
- // reshape (I) and the following power or multiply or both
- expected.push_back({9, 11});
- expected.push_back({9, 15});
- expected.push_back({9, 11, 15});
-
- // reshape (I) and the following power and multiply
- expected.push_back({9, 15, 17});
-
- // reshape (I) and everything after it to the max depth of 3
- expected.push_back({9, 11, 15, 17});
-
- // pairs of broadcasts
- expected.push_back({11, 24}); // multiply / multiply
- expected.push_back({15, 17}); // power / multiply
- expected.push_back({17, 18}); // multiply / add
- expected.push_back({18, 20}); // add / multiply
- expected.push_back({20, 21}); // multiply / tanh
- expected.push_back({21, 23}); // tanh / add
- expected.push_back({23, 24}); // add / multiply
-
- // triples of broadcasts
- expected.push_back({15, 17, 18}); // power / multiply / add
- expected.push_back({17, 18, 20}); // multiply / add / multiply
- expected.push_back({18, 20, 21}); // add / multiply / tanh
- expected.push_back({20, 21, 23}); // multiply / tanh / add
- expected.push_back({21, 23, 24}); // tanh / add / multiply
-
- auto expected_candidates = ExpectedCandidates(graph, spec, expected);
- AssertEqual(actual_candidates, expected_candidates);
- }
-}
-
-TEST(PartitionRule, Combine_TupleArg) {
- auto func = GPT2ExtractTuples();
- auto graph = DataflowGraph(func);
- auto spec = StandardSpec();
-
- {
- // Prime the system by picking out all 8 calls to non-opaque ops.
- auto sub_rule = OpCallByKindPartitionRule("op_call_by_kind");
- // Merge args of tuples of <= injective (I) fields into the call's group.
- Array<CombinerRule> combiner_rules;
- combiner_rules.push_back(TupleArgCombinerRule("tuple_arg"));
- // Build the overall partition rule.
- auto rule = CombinePartitionRule("combine_tuple_arg", std::move(sub_rule),
- std::move(combiner_rules), /*max_depth=*/3);
-
- auto actual_candidates = ActualCandidates(graph, func, spec, rule);
-
- // The original calls
- std::vector<std::vector<PostDfsIndex>> expected;
- expected.push_back({6});
- expected.push_back({8});
- expected.push_back({9});
- expected.push_back({10});
- expected.push_back({12});
- expected.push_back({13});
- expected.push_back({14});
- expected.push_back({16});
-
- // The concatenate((expand_dims(...), expand_dims(...)) is grouped.
- expected.push_back({10, 14, 15, 16});
-
- auto expected_candidates = ExpectedCandidates(graph, spec, expected);
- AssertEqual(actual_candidates, expected_candidates);
- }
-}
-
-TEST(PartitionRule, Combine_TupleProj) {
- auto func = GPT2ExtractTuples();
- auto graph = DataflowGraph(func);
- auto spec = StandardSpec();
-
- {
- // Prime the system by picking out all 8 calls to non-opaque ops.
- auto sub_rule = OpCallByKindPartitionRule("op_call_by_kind");
- // Merge projections from injective groups.
- Array<CombinerRule> combiner_rules;
- combiner_rules.push_back(TupleProjCombinerRule("tuple_proj"));
- // Build the overall partition rule.
- auto rule = CombinePartitionRule("combine_tuple_proj", std::move(sub_rule),
- std::move(combiner_rules), /*max_depth=*/3);
-
- auto actual_candidates = ActualCandidates(graph, func, spec, rule);
-
- // The original calls
- std::vector<std::vector<PostDfsIndex>> expected;
- expected.push_back({6});
- expected.push_back({8});
- expected.push_back({9});
- expected.push_back({10});
- expected.push_back({12});
- expected.push_back({13});
- expected.push_back({14});
- expected.push_back({16});
-
- // split / proj 1
- expected.push_back({6, 7});
- // split / proj 2
- expected.push_back({6, 11});
- // split and both projections
- expected.push_back({6, 7, 11});
-
- auto expected_candidates = ExpectedCandidates(graph, spec, expected);
- AssertEqual(actual_candidates, expected_candidates);
- }
-}
-
-TEST(PartitionRule, Combine_Constant) {
- auto func = GPT2ExtractOps();
- auto graph = DataflowGraph(func);
- auto spec = StandardSpec();
-
- {
- // Prime the system by picking out all 11 calls to non-opaque ops.
- auto sub_rule = OpCallByKindPartitionRule("op_call_by_kind");
- // Merge constant args into injective groups
- Array<CombinerRule> combiner_rules;
- combiner_rules.push_back(ConstantCombinerRule("constant"));
- // Build the overall partition rule.
- auto rule = CombinePartitionRule("combine_constant", std::move(sub_rule),
- std::move(combiner_rules), /*max_depth=*/3);
-
- auto actual_candidates = ActualCandidates(graph, func, spec, rule);
-
- // The original calls
- std::vector<std::vector<PostDfsIndex>> expected;
- expected.push_back({6});
- expected.push_back({8});
- expected.push_back({9});
- expected.push_back({11});
- expected.push_back({15});
- expected.push_back({17});
- expected.push_back({18});
- expected.push_back({20});
- expected.push_back({21});
- expected.push_back({23});
- expected.push_back({24});
-
- // Constant arg to nn.dense
- expected.push_back({5, 6});
-
- // Constant arg to add
- expected.push_back({7, 8});
-
- auto expected_candidates = ExpectedCandidates(graph, spec, expected);
- AssertEqual(actual_candidates, expected_candidates);
- }
-}
-
-TEST(PartitionRule, Combine_Mixed) {
- auto func = GPT2ExtractOps();
- auto graph = DataflowGraph(func);
- auto spec = StandardSpec();
-
- {
- // Prime the system by picking out all 11 calls to non-opaque ops.
- auto sub_rule = OpCallByKindPartitionRule("op_call_by_kind");
-
- // Mimic the FuseOps rules.
- Array<SimpleCombinerRule> simple_rules;
- simple_rules.push_back(ByKindSimpleCombinerRule(kOutEWiseFusable, kBroadcast));
- simple_rules.push_back(ByKindSimpleCombinerRule(kBroadcast, kCommReduce));
- simple_rules.push_back(ByKindSimpleCombinerRule(kInjective, kInjective));
- Array<CombinerRule> combiner_rules;
- combiner_rules.push_back(AllSimpleCombinerRule("all_simple", std::move(simple_rules)));
-
- // Merge constant args into injective groups
- combiner_rules.push_back(ConstantCombinerRule("constant"));
-
- // Build the overall partition rule.
- auto rule = CombinePartitionRule("combine_mixed", std::move(sub_rule),
- std::move(combiner_rules), /*max_depth=*/3);
-
- auto actual_candidates = ActualCandidates(graph, func, spec, rule);
-
- // The original calls
- std::vector<std::vector<PostDfsIndex>> expected;
- expected.push_back({6});
- expected.push_back({8});
- expected.push_back({9});
- expected.push_back({11});
- expected.push_back({15});
- expected.push_back({17});
- expected.push_back({18});
- expected.push_back({20});
- expected.push_back({21});
- expected.push_back({23});
- expected.push_back({24});
-
- // A -> B merging
- expected.push_back({6, 8});
- expected.push_back({9, 11});
- expected.push_back({9, 15});
- expected.push_back({9, 11, 15});
- expected.push_back({9, 15, 17});
- expected.push_back({9, 11, 15, 17});
- expected.push_back({11, 24});
- expected.push_back({15, 17});
- expected.push_back({17, 18});
- expected.push_back({18, 20});
- expected.push_back({20, 21});
- expected.push_back({21, 23});
- expected.push_back({23, 24});
- expected.push_back({15, 17, 18});
- expected.push_back({17, 18, 20});
- expected.push_back({18, 20, 21});
- expected.push_back({20, 21, 23});
- expected.push_back({21, 23, 24});
-
- // Constant args
- expected.push_back({5, 6});
- expected.push_back({7, 8});
-
- // B -> R
- expected.push_back({8, 9});
- expected.push_back({8, 9, 11});
- expected.push_back({8, 9, 15});
-
- // Constant's and A -> B
- expected.push_back({5, 6, 8});
- expected.push_back({5, 6, 7, 8});
-
- // Constants and B -> R
- expected.push_back({7, 8, 9});
- expected.push_back({7, 8, 9, 11});
- expected.push_back({7, 8, 9, 15});
-
- auto expected_candidates = ExpectedCandidates(graph, spec, expected);
- AssertEqual(actual_candidates, expected_candidates);
- }
-}
-
-TEST(PartitionRule, OnlyValid) {
- auto func = GPT2ExtractOps();
- auto graph = DataflowGraph(func);
- auto spec = StandardSpec();
-
- {
- // Prime the system by picking out all 11 calls to non-opaque ops.
- auto sub_rule = OpCallByKindPartitionRule("op_call_by_kind");
- // Combine all <= kOutEWiseFusable (A) actual_candidates (ie anything) with downstream
- // <= kBroadcast (B) actual_candidates (ie B or E).
- Array<SimpleCombinerRule> simple_rules;
- simple_rules.push_back(ByKindSimpleCombinerRule(/*upstream_kind=*/kOutEWiseFusable,
- /*downstream_kind=*/kBroadcast));
- Array<CombinerRule> combiner_rules;
- combiner_rules.push_back(AllSimpleCombinerRule("all_simple", std::move(simple_rules)));
- auto combine_rule = CombinePartitionRule("combine_by_kind_A_B", std::move(sub_rule),
- std::move(combiner_rules), /*max_depth=*/3);
- // Only allow up to depth 2, no taps and 1 exit.
- SubGraphConfig config;
- config.allow_taps = false;
- config.max_depth = 2;
- config.max_exits = 1;
-
- // Build the overall partition rule.
- auto rule = OnlyValidPartitionRule("only_valid", std::move(combine_rule), config);
-
- auto actual_candidates = ActualCandidates(graph, func, spec, rule);
-
- // The original calls.
- std::vector<std::vector<PostDfsIndex>> expected;
- expected.push_back({6});
- expected.push_back({8});
- expected.push_back({9});
- expected.push_back({11});
- expected.push_back({15});
- expected.push_back({17});
- expected.push_back({18});
- expected.push_back({20});
- expected.push_back({21});
- expected.push_back({23});
- expected.push_back({24});
-
- // nn.dense (A) and the following add (B)
- expected.push_back({6, 8});
-
- // pairs of broadcasts
- expected.push_back({11, 24}); // multiply / multiply
- expected.push_back({15, 17}); // power / multiply
- expected.push_back({17, 18}); // multiply / add
- expected.push_back({18, 20}); // add / multiply
- expected.push_back({20, 21}); // multiply / tanh
- expected.push_back({21, 23}); // tanh / add
- expected.push_back({23, 24}); // add / multiply
-
- // The following candidates are filtered out because they have 2 or 3 exits:
- // {9, 11}, {9, 15}, {9,11,15}, {9,15,17}, {15,17,18}, {17,18,20},
- // {18,20,21}, {20,21,23}, {21,23,24}, {9,11,15,17}
-
- auto expected_candidates = ExpectedCandidates(graph, spec, expected);
- AssertEqual(actual_candidates, expected_candidates);
- }
-}
-
-TEST(PartitionRule, Host) {
- auto func = GPT2ExtractTuples();
- auto graph = DataflowGraph(func);
- auto spec = StandardSpec();
-
- {
- auto rule = HostPartitionRule("host");
-
- auto actual_candidates = ActualCandidates(graph, func, spec, rule);
-
- std::vector<std::vector<PostDfsIndex>> expected;
-
- // Function arg %x
- expected.push_back({0});
- // Operators
- expected.push_back({1}); // concatenate
- expected.push_back({2}); // expand_dims
- expected.push_back({3}); // transpose
- expected.push_back({4}); // reshape
- expected.push_back({5}); // split
- // Tuple projection
- expected.push_back({7});
- expected.push_back({11});
- // Tuple construction
- expected.push_back({15});
- // The overall @main function
- expected.push_back({17});
-
- auto expected_candidates = ExpectedCandidates(graph, spec, expected);
- AssertEqual(actual_candidates, expected_candidates);
- }
-}
-
-} // namespace
-} // namespace collage
-} // namespace relay
-} // namespace tvm
diff --git a/tests/cpp/relay/df_pattern_rewrite_test.cc b/tests/cpp/relay/df_pattern_rewrite_test.cc
deleted file mode 100644
index 374887c12a..0000000000
--- a/tests/cpp/relay/df_pattern_rewrite_test.cc
+++ /dev/null
@@ -1,100 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include <gtest/gtest.h>
-#include <tvm/relay/attrs/transform.h>
-#include <tvm/relay/dataflow_matcher.h>
-#include <tvm/relay/dataflow_pattern.h>
-#include <tvm/relay/function.h>
-#include <tvm/relay/parser.h>
-
-#include "../../../src/relay/transforms/simplify_expr.h"
-
-namespace tvm {
-namespace relay {
-namespace {
-
-// Demonstrates rewriting a deeply nested sub-graph with specific
-// attributes on the inner-most operator call.
-class TestRewriter : public DFPatternRewrite {
- public:
- TestRewriter() {
- x_ = IsWildcard();
- const1_ = IsWildcard();
- const2_ = IsWildcard();
- const3_ = IsWildcard();
- const4_ = IsWildcard();
-
- auto biasadd = IsOp("nn.bias_add");
- auto relu = IsOp("nn.relu");
- auto conv2d = IsOp("nn.conv2d");
-
- Map<String, ObjectRef> attrs;
- attrs.Set("groups", Integer(304));
- auto maybedepthwise = conv2d({x_, const1_}).HasAttr(attrs);
-
- pattern_ =
- relu({biasadd({conv2d({relu({biasadd({maybedepthwise, const2_})}), const3_}), const4_})});
- }
-
- Expr Callback(const Expr& pre, const Expr& post,
- const Map<DFPattern, Array<Expr>>& node_map) const override {
- LOG(INFO) << "depthwise conv2d detected!";
- auto attrs = runtime::make_object<InitOpAttrs>();
- attrs->shape = Array<Integer>({Integer(1), Integer(256), Integer(128), Integer(128)});
- attrs->dtype = DataType::Float(32);
- return Call(Op::Get("zeros"), {}, Attrs(attrs));
- }
-
- DFPattern x_, const1_, const2_, const3_, const4_;
-};
-
-TEST(DFPatternRewrite, DeeplyNestedWithCallAttributes) {
- constexpr const char* kModel = R"(
- #[version = "0.0.5"]
- def @main(%data : Tensor[(1, 304, 128, 128), float32],
- %weight1 : Tensor[(304, 1, 3, 3), float32],
- %bias1 : Tensor[(304), float32],
- %weight2 : Tensor[(256, 304, 1, 1), float32],
- %bias2 : Tensor[(256), float32]) -> Tensor[(1, 256, 128, 128), float32] {
- %0 = nn.conv2d(%data, %weight1, padding=[1, 1, 1, 1], groups=304, channels=304, kernel_size=[3, 3]);
- %1 = nn.bias_add(%0, %bias1);
- %2 = nn.relu(%1);
- %3 = nn.conv2d(%2, %weight2, padding=[0, 0, 0, 0], channels=256, kernel_size=[1, 1]);
- %4 = nn.bias_add(%3, %bias2);
- nn.relu(%4)
- }
- )";
-
- IRModule module = ParseModule("string", kModel);
- DFPatternRewriteComposer composer;
- composer.AddRewrite<TestRewriter>();
- Function in_function = Downcast<Function>(module->Lookup("main"));
- LOG(INFO) << "input function:\n" << PrettyPrint(in_function);
- Function out_function =
- Downcast<Function>(RewritePatterns(composer.MakeCallbacks(), in_function, module));
- LOG(INFO) << "output function:\n" << PrettyPrint(out_function);
- const auto* call_node = out_function->body.as<CallNode>();
- ASSERT_TRUE(call_node != nullptr);
- ASSERT_TRUE(call_node->op == Op::Get("zeros"));
-}
-
-} // namespace
-} // namespace relay
-} // namespace tvm
diff --git a/tests/cpp/relay/ir/indexed_graph_test.cc b/tests/cpp/relay/ir/indexed_graph_test.cc
deleted file mode 100644
index 486d027fbc..0000000000
--- a/tests/cpp/relay/ir/indexed_graph_test.cc
+++ /dev/null
@@ -1,205 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include "../../../src/relay/ir/indexed_graph.h"
-
-#include <gtest/gtest.h>
-#include <tvm/relay/expr.h>
-#include <tvm/relay/function.h>
-#include <tvm/relay/parser.h>
-
-namespace tvm {
-namespace relay {
-namespace {
-
-// A module stolen from onnx/test_forward.py::test_loop which combines functions, recursion,
-// control flow, tuples as well as the usual operator calls.
-// We include the known post-dfs indexes in comments to help write the tests.
-IRModule TestRecursiveIRModule() {
- Device device = {kDLCPU, 0};
- Constant const0(runtime::NDArray::Empty(ShapeTuple({1}), DataType::Int(64), device));
- Constant const1(runtime::NDArray::Empty(ShapeTuple({0, 1}), DataType::Float(32), device));
- Map<String, Array<ObjectRef>> metadata;
- metadata.Set("relay.Constant", Array<ObjectRef>({const0, const1}));
- constexpr const char* kModel = R"(
- #[version = "0.0.5"]
- def @main(%trip_count: int64, // 0
- %cond: bool, // 1
- %y: Tensor[(1), float32]) // 2
- -> (Tensor[(1), float32], Tensor[(?, ?), float32]) {
- %17 = (
- let %while_loop = fn (%iter_count: int64, // 3
- %max_count: int64, // 4
- %cond_in: bool, // 5
- %y_in: Tensor[(1), float32], // 6
- %scan_out: Tensor[(?, ?), float32]) // 7
- -> (int64, int64, bool, Tensor[(1), float32], Tensor[(?, ?), float32]) {
- %0 = equal(%cond_in, True); // 11
- %1 = less(%iter_count, %max_count); // 13
- %2 = logical_and(%0, %1); // 14
- if (%2) {
- %3 = cast(%iter_count, dtype="float32"); // 20
- %4 = add(%y_in, %3); // 21
- %5 = less(%4, 5f); // 23
- %6 = squeeze(%5); // 24
- %7 = reshape(%iter_count, newshape=[1]); // 29
- %8 = (%7, meta[relay.Constant][0]); // 31
- %9 = concatenate(%8); // 32
- %10 = copy(%4); // 36
- %11 = dyn.broadcast_to(%scan_out, %9, shape=None); // 33
- %12 = expand_dims(%10, axis=0); // 37
- %13 = (%11, %12); // 38
- %14 = add(%iter_count, 1i64); // 17
- %15 = cast(%6, dtype="bool"); // 25
- %16 = concatenate(%13); // 39
- %while_loop(%14, %max_count, %15, %4, %16) // 40
- } else {
- (%iter_count, %max_count, %cond_in, %y_in, %scan_out) // 41
- } // 42
- }; // 43
- %while_loop // 44
- ); // 45
- %18 = %17(0i64, %trip_count, %cond, %y, meta[relay.Constant][1]); // 48
- %19 = %18.3; // 49
- %20 = %18.4; // 50
- (%19, %20) // 51
- } // 52
- )";
- return ParseModule("string", kModel, /*init_module=*/{}, metadata);
-}
-
-TEST(IndexedGraph, RecursiveExprRegression) {
- IRModule ir_mod = TestRecursiveIRModule();
- auto main = Downcast<Function>(ir_mod->Lookup("main"));
- auto graph = CreateIndexedGraph(main);
- graph->CheckValid();
-
- {
- // Dataflow node properties for %4
- auto node = graph->index_to_node(21);
- const auto* call_node = node->ref().as<CallNode>();
- ASSERT_NE(call_node, nullptr);
- const auto* op_node = call_node->op.as<OpNode>();
- ASSERT_NE(op_node, nullptr);
- ASSERT_EQ(op_node->name, "add");
-
- // 3 inputs (the op itself is an input)
- ASSERT_EQ(node->inputs_.size(), 3);
- ASSERT_EQ(node->inputs_[0]->index_, 15); // the add op
- ASSERT_EQ(node->inputs_[1]->index_, 6); // %y_in
- ASSERT_EQ(node->inputs_[2]->index_, 20); // %3
-
- // 3 outputs
- ASSERT_EQ(node->outputs_.size(), 3);
- ASSERT_EQ(node->outputs_[0]->index_, 23); // %5
- ASSERT_EQ(node->outputs_[1]->index_, 36); // %10
- ASSERT_EQ(node->outputs_[2]->index_, 40); // recursive %while_loop call
-
- // In the 'if' basic block
- ASSERT_EQ(node->basic_block_->index_, 42);
-
- // Dominator 'parent' is recursive call
- ASSERT_EQ(node->dominator_parent_->index_, 40);
-
- // One dominator child from %3
- ASSERT_EQ(node->dominator_children_.size(), 1);
- ASSERT_EQ(node->dominator_children_[0]->index_, 20);
- }
-
- {
- // The recursive call to %while_loop does not depend on %while_loop
- auto node = graph->index_to_node(40);
- const auto* call_node = node->ref().as<CallNode>();
- ASSERT_NE(call_node, nullptr);
- const auto* var_node = call_node->op.as<VarNode>();
- ASSERT_NE(var_node, nullptr);
- ASSERT_EQ(var_node->name_hint(), "while_loop");
-
- ASSERT_EQ(node->inputs_.size(), 5);
- ASSERT_EQ(node->inputs_[0]->index_, 17); // %14
- ASSERT_EQ(node->inputs_[1]->index_, 4); // %max_count
- ASSERT_EQ(node->inputs_[2]->index_, 25); // %15
- ASSERT_EQ(node->inputs_[3]->index_, 21); // %4
- ASSERT_EQ(node->inputs_[4]->index_, 39); // %16
- }
-
- {
- // Downstream nodes of %18
- auto node = graph->index_to_node(48);
- std::unordered_set<const IndexedGraph<Expr>::Node*> downstreams;
- node->AccumulateDownstreamNodes(&downstreams);
- ASSERT_EQ(downstreams.size(), 4);
- for (const auto* downstream : downstreams) {
- ASSERT_TRUE(downstream->index_ >= 49 && downstream->index_ <= 52);
- }
- }
-
- {
- // Dominates relation for %4
- auto upstream = graph->index_to_node(21);
- // Path 1: 21->23->24->25->40
- // Path 2: 21->36->37->38->39->40
- // Then 40->43
- auto downstream = graph->index_to_node(43);
- ASSERT_TRUE(downstream->Dominates(upstream));
- }
-}
-
-// A module with unused let-bound function. The 'add' operator should have no dominator
-// since it is used both in the unused function and in the main body.
-IRModule TestUnusedLetBoundIRModule() {
- constexpr const char* kModel = R"(
- #[version = "0.0.5"]
- def @main(%x: int64) -> int64 { // 0
- let %f = fn ( // 5
- %y: int64 // 1
- ) {
- add(%x, %y) // 3
- };
- if (less(%x, 5i64)) {
- add(%x, 3i64) // 10
- } else {
- %x
- }
- }
- )";
- return ParseModule("string", kModel);
-}
-
-TEST(IndexedGraph, UnusedLetVars) {
- IRModule ir_mod = TestUnusedLetBoundIRModule();
- auto main = Downcast<Function>(ir_mod->Lookup("main"));
- auto graph = CreateIndexedGraph(main);
- graph->CheckValid();
-
- {
- auto node = graph->index_to_node(2);
- const auto* op_node = node->ref().as<OpNode>();
- ICHECK(op_node);
- ICHECK_EQ(op_node->name, "add");
- ICHECK_EQ(node->outputs_.size(), 2);
- ICHECK_EQ(node->outputs_[0]->index_, 3);
- ICHECK_EQ(node->outputs_[1]->index_, 10);
- ICHECK(node->dominator_parent_ == nullptr);
- }
-}
-
-} // namespace
-} // namespace relay
-} // namespace tvm
diff --git a/tests/cpp/relay/op/memory/on_device_test.cc b/tests/cpp/relay/op/memory/on_device_test.cc
deleted file mode 100644
index 6f0a0b0d8b..0000000000
--- a/tests/cpp/relay/op/memory/on_device_test.cc
+++ /dev/null
@@ -1,89 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include "../../../../../src/relay/op/memory/on_device.h"
-
-#include <gtest/gtest.h>
-
-#include <string>
-
-namespace tvm {
-namespace relay {
-
-TEST(OnDeviceOp, Name) { EXPECT_EQ(OnDeviceOp()->name, "on_device"); }
-
-TEST(OnDevice, Default) {
- Var body("x", {});
- VirtualDevice virtual_device = VirtualDevice::ForDeviceType(kDLCPU, 3);
- Call call = OnDevice(body, virtual_device);
- EXPECT_EQ(call->op, OnDeviceOp());
- EXPECT_EQ(call->args.size(), 1);
- EXPECT_EQ(call->args[0], body);
- const auto* attrs = call->attrs.as<OnDeviceAttrs>();
- ASSERT_TRUE(attrs != nullptr);
- EXPECT_EQ(attrs->virtual_device, virtual_device);
- EXPECT_FALSE(attrs->constrain_result);
- EXPECT_TRUE(attrs->constrain_body);
-}
-
-TEST(OnDevice, Fixed) {
- Var body("x", {});
- VirtualDevice virtual_device = VirtualDevice::ForDeviceType(kDLCPU, 3);
- Call call = OnDevice(body, virtual_device, /*constrain_result=*/true);
- const auto* attrs = call->attrs.as<OnDeviceAttrs>();
- ASSERT_TRUE(attrs != nullptr);
- EXPECT_TRUE(attrs->constrain_result);
- EXPECT_TRUE(attrs->constrain_body);
-}
-
-TEST(OnDevice, Free) {
- Var body("x", {});
- VirtualDevice virtual_device = VirtualDevice::ForDeviceType(kDLCPU, 3);
- Call call = OnDevice(body, virtual_device, /*constrain_result=*/false, /*constrain_body=*/false);
- const auto* attrs = call->attrs.as<OnDeviceAttrs>();
- ASSERT_TRUE(attrs != nullptr);
- EXPECT_FALSE(attrs->constrain_result);
- EXPECT_FALSE(attrs->constrain_body);
-}
-
-TEST(GetOnDeviceProps, Correct) {
- Var body("x", {});
- VirtualDevice virtual_device = VirtualDevice::ForDeviceType(kDLCPU, 3);
- Call call = OnDevice(body, virtual_device, /*constrain_result=*/true, /*constrain_body=*/false);
- OnDeviceProps props = GetOnDeviceProps(call);
- ASSERT_TRUE(props.body.defined());
- ASSERT_EQ(props.virtual_device, virtual_device);
- ASSERT_TRUE(props.constrain_result);
- ASSERT_FALSE(props.constrain_body);
-}
-
-TEST(MaybeOnDevice, Wrapped) {
- VirtualDevice virtual_device = VirtualDevice::ForDeviceType(kDLCPU, 3);
- Var body("x", {});
- Call inner = OnDevice(body, virtual_device);
- Call outer = OnDevice(inner, virtual_device);
- OnDeviceProps props = GetOnDeviceProps(outer);
- ASSERT_TRUE(props.body.defined());
- ASSERT_EQ(props.virtual_device, virtual_device);
- ASSERT_FALSE(props.constrain_result);
- ASSERT_TRUE(props.constrain_body);
-}
-
-} // namespace relay
-} // namespace tvm
diff --git a/tests/cpp/relay/transforms/device_domains_test.cc b/tests/cpp/relay/transforms/device_domains_test.cc
deleted file mode 100644
index 47e303996b..0000000000
--- a/tests/cpp/relay/transforms/device_domains_test.cc
+++ /dev/null
@@ -1,76 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-/*
- * Just a smoke test for the device planner's unification domain, mostly to tease out how we'd
- * like to organize our cpp unit tests for functionality that's not obviously a Pass or should
- * be exposed via FFI.
- */
-
-// TODO(mbs): Revisit cpp unit test layout or setup include dir at root of src/
-#include "../../../../src/relay/transforms/device_domains.h"
-
-#include <gtest/gtest.h>
-#include <tvm/relay/parser.h>
-#include <tvm/relay/transform.h>
-
-namespace tvm {
-namespace relay {
-namespace transform {
-namespace {
-
-IRModule TestModule() {
- return InferType()(ParseModule("test", R"(
- #[version = "0.0.5"]
- def @f(%x : Tensor[(3, 7), float32], %y : Tensor[(3, 7), float32]) {
- add(%x, %y)
- }
- )"));
-}
-
-TEST(DeviceDomains, SmokeTest) {
- VirtualDevice cpu = VirtualDevice::ForDeviceType(kDLCPU);
- VirtualDevice cuda = VirtualDevice::ForDeviceType(kDLCUDA);
- transform::PassContext ctxt = transform::PassContext::Create();
- CompilationConfig config(ctxt, {Target("llvm"), Target("cuda")});
- DeviceDomains domains(config);
- IRModule mod = TestModule();
- Function f = Downcast<Function>(mod->Lookup("f"));
-
- DeviceDomainPtr actual_add_domain = domains.DomainForCallee(Downcast<Call>(f->body));
- DeviceDomainPtr x_domain = domains.DomainFor(f->params[0]);
- DeviceDomainPtr y_domain = domains.DomainFor(f->params[1]);
- DeviceDomainPtr result_domain = domains.Free(f->ret_type);
- std::vector<DeviceDomainPtr> arg_and_results;
- arg_and_results.push_back(x_domain);
- arg_and_results.push_back(y_domain);
- arg_and_results.push_back(result_domain);
- DeviceDomainPtr implied_add_domain = domains.MakeHigherOrderDomain(std::move(arg_and_results));
- EXPECT_FALSE(domains.UnifyOrNull(actual_add_domain, implied_add_domain) == nullptr);
- EXPECT_FALSE(domains.UnifyOrNull(x_domain, domains.ForVirtualDevice(f->params[0]->checked_type(),
- cuda)) == nullptr);
-
- EXPECT_EQ(domains.ResultVirtualDevice(y_domain), config->CanonicalVirtualDevice(cuda));
- EXPECT_EQ(domains.ResultVirtualDevice(result_domain), config->CanonicalVirtualDevice(cuda));
-}
-
-} // namespace
-} // namespace transform
-} // namespace relay
-} // namespace tvm
diff --git a/tests/cpp/relay/with_fields_test.cc b/tests/cpp/relay/with_fields_test.cc
deleted file mode 100644
index 6114fa97a9..0000000000
--- a/tests/cpp/relay/with_fields_test.cc
+++ /dev/null
@@ -1,215 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and lixmitations
- * under the License.
- */
-
-/*!
- * \brief Proof-of-concept unit tests for the family of WithFields helpers.
- * Only Call, GlobalVar and Constant are currently tested.
- */
-
-#include <gtest/gtest.h>
-#include <tvm/relay/adt.h>
-#include <tvm/relay/expr.h>
-#include <tvm/relay/function.h>
-#include <tvm/relay/parser.h>
-
-namespace tvm {
-namespace relay {
-namespace {
-
-IRModule TestIRModule() {
- return ParseModule("string",
- R"(
- #[version = "0.0.5"]
- def @main(%data : Tensor[(1, 304, 128, 128), float32],
- %weight1 : Tensor[(304, 1, 3, 3), float32],
- %bias1 : Tensor[(304), float32],
- %weight2 : Tensor[(256, 304, 1, 1), float32],
- %bias2 : Tensor[(256), float32]) -> Tensor[(1, 256, 128, 128), float32] {
- %0 = nn.conv2d(%data, %weight1, padding=[1, 1, 1, 1], groups=304, channels=304, kernel_size=[3, 3]);
- %1 = nn.bias_add(%0, %bias1);
- %2 = nn.relu(%1);
- %3 = nn.conv2d(%2, %weight2, padding=[0, 0, 0, 0], channels=256, kernel_size=[1, 1]);
- %4 = nn.bias_add(%3, %bias2);
- nn.relu(%4)
- }
- )");
-}
-
-Function TestFunction() { return Downcast<Function>(TestIRModule()->Lookup("main")); }
-Call TestCall() { return Downcast<Call>(TestFunction()->body); }
-GlobalVar TestGlobalVar() { return TestIRModule()->GetGlobalVar("main"); }
-VirtualDevice TestVirtualDevice() { return VirtualDevice::ForDevice({kDLCUDA, 3}); }
-Span TestSpan() { return Span(SourceName::Get("foo"), 3, 4, 6, 42); }
-Constant TestConstant() {
- return Constant(runtime::NDArray::Empty({}, DataType::Int(32), {kDLCPU, 0}));
-}
-
-//
-// Call
-//
-
-TEST(WithFields, Call_Noop) {
- Call call = TestCall();
- Call result = WithFields(call);
- ASSERT_TRUE(result.same_as(call));
-}
-
-TEST(WithFields, Call_Op) {
- Call call = TestCall();
- Op new_op = Op::Get("tanh");
- Call result = WithFields(call, new_op);
- ASSERT_FALSE(result.same_as(call));
- ASSERT_FALSE(call->op.same_as(new_op));
- ASSERT_TRUE(result->op.same_as(new_op));
-}
-
-TEST(WithFields, Call_Args) {
- Call call = TestCall();
- Array<Expr> new_args = {Tuple(Array<Expr>())};
- Call result = WithFields(call, /*opt_op=*/{}, new_args);
- ASSERT_FALSE(result.same_as(call));
- ASSERT_FALSE(call->args.same_as(new_args));
- ASSERT_TRUE(result->args.same_as(new_args));
-}
-
-TEST(WithFields, Call_Attrs) {
- Call call = TestCall();
- Attrs new_attrs = DictAttrs(Map<String, ObjectRef>());
- Call result = WithFields(call, /*opt_op=*/{}, /*opt_args=*/{}, new_attrs);
- ASSERT_FALSE(result.same_as(call));
- ASSERT_FALSE(call->attrs.same_as(new_attrs));
- ASSERT_TRUE(result->attrs.same_as(new_attrs));
-}
-
-TEST(WithFields, Call_TypeArgs) {
- Call call = TestCall();
- Array<Type> new_type_args;
- Call result = WithFields(call, /*opt_op=*/{}, /*opt_args=*/{}, /*opt_attrs=*/{}, new_type_args);
- ASSERT_FALSE(result.same_as(call));
- ASSERT_FALSE(call->type_args.same_as(new_type_args));
- ASSERT_TRUE(result->type_args.same_as(new_type_args));
-}
-
-TEST(WithFields, Call_VirtualDevice) {
- Call call = TestCall();
- VirtualDevice new_virtual_device = TestVirtualDevice();
- Call result = WithFields(call, /*opt_op=*/{}, /*opt_args=*/{}, /*opt_attrs=*/{},
- /*opt_type_args=*/{}, new_virtual_device);
- ASSERT_FALSE(result.same_as(call));
- ASSERT_FALSE(call->virtual_device().same_as(new_virtual_device));
- ASSERT_TRUE(result->virtual_device().same_as(new_virtual_device));
-}
-
-TEST(WithFields, Call_Span) {
- Call call = TestCall();
- Span new_span = TestSpan();
- Call result = WithFields(call, /*opt_op=*/{}, /*opt_args=*/{}, /*opt_attrs=*/{},
- /*opt_type_args=*/{}, /*opt_virtual_device=*/{}, new_span);
- ASSERT_FALSE(result.same_as(call));
- ASSERT_FALSE(call->span.same_as(new_span));
- ASSERT_TRUE(result->span.same_as(new_span));
-}
-
-//
-// GlobalVar
-//
-
-TEST(WithFields, GlobalVar_Noop) {
- GlobalVar gv = TestGlobalVar();
- GlobalVar result = WithFields(gv);
- ASSERT_TRUE(result.same_as(gv));
-}
-
-TEST(WithFields, GlobalVar_Name) {
- GlobalVar gv = TestGlobalVar();
- String new_name("foo");
- GlobalVar result = WithFields(gv, new_name);
- ASSERT_FALSE(result.same_as(gv));
- ASSERT_FALSE(gv->name_hint.same_as(new_name));
- ASSERT_TRUE(result->name_hint.same_as(new_name));
-}
-
-TEST(WithFields, GlobalVar_Type) {
- GlobalVar gv = TestGlobalVar();
- Type new_type = TupleType(Array<Type>());
- GlobalVar result = WithFields(gv, /*opt_name_hint=*/{}, new_type);
- ASSERT_FALSE(result.same_as(gv));
- ASSERT_FALSE(gv->checked_type().same_as(new_type));
- ASSERT_TRUE(result->checked_type().same_as(new_type));
-}
-
-TEST(WithFields, GlobalVar_VirtualDevice) {
- GlobalVar gv = TestGlobalVar();
- VirtualDevice new_virtual_device = TestVirtualDevice();
- GlobalVar result = WithFields(gv, /*opt_name_hint=*/{}, /*opt_type=*/{}, new_virtual_device);
- ASSERT_FALSE(result.same_as(gv));
- ASSERT_FALSE(gv->virtual_device().same_as(new_virtual_device));
- ASSERT_TRUE(result->virtual_device().same_as(new_virtual_device));
-}
-
-TEST(WithFields, GlobalVar_Span) {
- GlobalVar gv = TestGlobalVar();
- Span new_span = TestSpan();
- GlobalVar result =
- WithFields(gv, /*opt_name_hint=*/{}, /*opt_type=*/{}, /*opt_virtual_device=*/{}, new_span);
- ASSERT_FALSE(result.same_as(gv));
- ASSERT_FALSE(gv->span.same_as(new_span));
- ASSERT_TRUE(result->span.same_as(new_span));
-}
-
-//
-// Constant
-//
-
-TEST(WithFields, Constant_Noop) {
- Constant constant = TestConstant();
- Constant result = WithFields(constant);
- ASSERT_TRUE(result.same_as(constant));
-}
-
-TEST(WithFields, Constant_Data) {
- Constant constant = TestConstant();
- runtime::NDArray new_data = runtime::NDArray::Empty({}, DataType::Float(32), {kDLCPU, 0});
- Constant result = WithFields(constant, new_data);
- ASSERT_FALSE(result.same_as(constant));
- ASSERT_FALSE(constant->data.same_as(new_data));
- ASSERT_TRUE(result->data.same_as(new_data));
-}
-
-TEST(WithFields, Constant_VirtualDevice) {
- Constant constant = TestConstant();
- VirtualDevice new_virtual_device = TestVirtualDevice();
- Constant result = WithFields(constant, /*opt_data=*/{}, new_virtual_device);
- ASSERT_FALSE(result.same_as(constant));
- ASSERT_FALSE(constant->virtual_device().same_as(new_virtual_device));
- ASSERT_TRUE(result->virtual_device().same_as(new_virtual_device));
-}
-
-TEST(WithFields, Constant_Span) {
- Constant constant = TestConstant();
- Span new_span = TestSpan();
- Constant result = WithFields(constant, /*opt_data=*/{}, /*opt_virtual_device=*/{}, new_span);
- ASSERT_FALSE(result.same_as(constant));
- ASSERT_FALSE(constant->span.same_as(new_span));
- ASSERT_TRUE(result->span.same_as(new_span));
-}
-
-} // namespace
-} // namespace relay
-} // namespace tvm
diff --git a/tests/cpp/relay_build_module_test.cc b/tests/cpp/relay_build_module_test.cc
deleted file mode 100644
index 103989fc77..0000000000
--- a/tests/cpp/relay_build_module_test.cc
+++ /dev/null
@@ -1,193 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include <gtest/gtest.h>
-#include <tvm/driver/driver_api.h>
-#include <tvm/ir/memory_pools.h>
-#include <tvm/ir/module.h>
-#include <tvm/relay/analysis.h>
-#include <tvm/relay/executor.h>
-#include <tvm/relay/expr.h>
-#include <tvm/relay/op_attr_types.h>
-#include <tvm/relay/op_strategy.h>
-#include <tvm/relay/runtime.h>
-#include <tvm/relay/transform.h>
-#include <tvm/relay/type.h>
-#include <tvm/runtime/executor_info.h>
-#include <tvm/runtime/module.h>
-#include <tvm/runtime/packed_func.h>
-#include <tvm/runtime/registry.h>
-#include <tvm/te/operation.h>
-#include <tvm/topi/broadcast.h>
-#include <tvm/topi/generic/injective.h>
-
-using namespace tvm;
-using namespace tvm::relay;
-
-TVM_REGISTER_GLOBAL("test.strategy")
- .set_body_typed([](const Attrs& attrs, const Array<te::Tensor>& inputs, const Type& out_type,
- const Target& target) {
- FTVMCompute fcompute = [](const Attrs& attrs, const Array<te::Tensor>& inputs,
- const Type& out_type) -> Array<te::Tensor> {
- ICHECK_EQ(inputs.size(), 2U);
- return {topi::add(inputs[0], inputs[1])};
- };
- FTVMSchedule fschedule = [](const Attrs& attrs, const Array<te::Tensor>& outs,
- const Target& target) {
- With<Target> target_scope(target);
- return topi::generic::schedule_injective(target, outs);
- };
-
- auto n = make_object<OpStrategyNode>();
- auto strategy = tvm::relay::OpStrategy(std::move(n));
- strategy.AddImplementation(fcompute, fschedule, "test.strategy", 10);
- return strategy;
- });
-
-TVM_REGISTER_GLOBAL("relay.backend.lower_call")
- .set_body_typed([](const relay::Call& call, const Array<te::Tensor>& inputs,
- const Target& target) {
- static auto fstrategy = Op::GetAttrMap<relay::FTVMStrategy>("FTVMStrategy");
- Op op = Downcast<Op>(call->op);
- auto out_type = call->checked_type();
- OpStrategy strategy = fstrategy[op](call->attrs, inputs, out_type, target);
- auto impl = strategy->specializations[0]->implementations[0];
- auto outs = impl.Compute(call->attrs, inputs, out_type);
- auto f = tvm::runtime::Registry::Get("relay.backend._make_LoweredOutput");
- if (!f) {
- LOG(FATAL) << "relay.backend._make_LoweredOutput is not registered";
- }
- return (*f)(outs, impl);
- });
-
-TEST(Relay, BuildModule) {
- auto tensor_type = relay::TensorType({2, 3}, DataType::Float(32));
- auto a = relay::Var("a", tensor_type);
- auto b = relay::Var("b", tensor_type);
- auto add_op = relay::Op::Get("add");
- auto x = relay::Call(add_op, {a, b}, tvm::Attrs(), {});
- auto c = relay::Var("c", tensor_type);
- auto y = relay::Call(add_op, {x, c}, tvm::Attrs(), {});
- auto func = relay::Function(relay::FreeVars(y), y, relay::Type(), {});
- auto A = tvm::runtime::NDArray::Empty({2, 3}, {kDLFloat, 32, 1}, {kDLCPU, 0});
- auto B = tvm::runtime::NDArray::Empty({2, 3}, {kDLFloat, 32, 1}, {kDLCPU, 0});
- auto C = tvm::runtime::NDArray::Empty({2, 3}, {kDLFloat, 32, 1}, {kDLCPU, 0});
-
- auto pA = static_cast<float*>(A->data);
- auto pB = static_cast<float*>(B->data);
- auto pC = static_cast<float*>(C->data);
-
- for (int i = 0; i < 6; ++i) {
- pA[i] = i;
- pB[i] = i + 1;
- pC[i] = i + 2;
- }
- // get schedule
- auto reg = tvm::runtime::Registry::Get("ir.RegisterOpAttr");
- if (!reg) {
- LOG(FATAL) << "no _Register";
- }
- auto reset = tvm::runtime::Registry::Get("ir.OpResetAttr");
- if (!reset) {
- LOG(FATAL) << "Reset is not defined.";
- }
- auto fs = tvm::runtime::Registry::Get("test.strategy");
- if (!fs) {
- LOG(FATAL) << "No test_strategy registered.";
- }
- auto fgeneric = GenericFunc::Get("test.strategy_generic").set_default(*fs, true);
- (*reset)(add_op, "FTVMStrategy");
- (*reg)("add", "FTVMStrategy", fgeneric, 10);
- Array<Integer> dep;
- dep.push_back(0);
- (*reset)(add_op, "TShapeDataDependent");
- (*reg)("add", "TShapeDataDependent", dep, 10);
- // build
- auto pfb = tvm::runtime::Registry::Get("relay.build_module._BuildModule");
- tvm::runtime::Module build_mod = (*pfb)();
- auto build_f = build_mod.GetFunction("build", false);
- auto json_f = build_mod.GetFunction("get_graph_json", false);
- auto mod_f = build_mod.GetFunction("get_module", false);
- Target llvm_tgt = Target("llvm");
- Array<Target> targets = {llvm_tgt};
- auto relay_mod = tvm::IRModule::FromExpr(func);
- ICHECK(relay_mod.defined()) << "Module must be defined";
- build_f(relay_mod, targets, llvm_tgt, Executor::Create("graph"), Runtime::Create("cpp"),
- WorkspaceMemoryPools(), ConstantMemoryPools(), "");
- std::string json = json_f();
- tvm::runtime::Module mod = mod_f();
- // run
- auto dev = A->device;
- auto pfr = tvm::runtime::Registry::Get("tvm.graph_executor.create");
- ICHECK(mod.defined()) << "Module must be defined";
- tvm::runtime::Module run_mod =
- (*pfr)(json, mod, static_cast<int>(dev.device_type), dev.device_id);
- auto set_input_f = run_mod.GetFunction("set_input_zero_copy", false);
- auto run_f = run_mod.GetFunction("run", false);
- auto get_output_f = run_mod.GetFunction("get_output", false);
- set_input_f("a", const_cast<DLTensor*>(A.operator->()));
- set_input_f("b", const_cast<DLTensor*>(B.operator->()));
- set_input_f("c", const_cast<DLTensor*>(C.operator->()));
- run_f();
- tvm::runtime::NDArray Y = get_output_f(0);
- auto pY = static_cast<float*>(Y->data);
- for (int i = 0; i < 6; ++i) {
- ICHECK_LT(fabs(pY[i] - (i + (i + 1) + (i + 2))), 1e-4);
- }
- // mutate the input a bit and run it again
- for (int i = 0; i < 6; ++i) {
- pB[i] = i + 3;
- }
- run_f();
- tvm::runtime::NDArray Y2 = get_output_f(0);
- auto pY2 = static_cast<float*>(Y2->data);
- for (int i = 0; i < 6; ++i) {
- ICHECK_LT(fabs(pY2[i] - (i + (i + 3) + (i + 2))), 1e-4);
- }
- // attach a different input and run it again
- auto C2 = tvm::runtime::NDArray::Empty({2, 3}, {kDLFloat, 32, 1}, {kDLCPU, 0});
- auto pC2 = static_cast<float*>(C2->data);
- for (int i = 0; i < 6; ++i) {
- pC2[i] = i + 4;
- }
- set_input_f("c", const_cast<DLTensor*>(C2.operator->()));
- run_f();
- tvm::runtime::NDArray Y3 = get_output_f(0);
- auto pY3 = static_cast<float*>(Y3->data);
- for (int i = 0; i < 6; ++i) {
- ICHECK_LT(fabs(pY3[i] - (i + (i + 3) + (i + 4))), 1e-4);
- }
-}
-
-TEST(Relay, GetExprRefCount) {
- auto tensor_type = relay::TensorType({2, 3}, DataType::Float(32));
- auto a = relay::Var("a", tensor_type);
- auto add_op = relay::Op::Get("add");
- auto relu_op = relay::Op::Get("nn.relu");
- auto x = relay::Call(relu_op, {a}, tvm::Attrs(), {});
- auto y = relay::Call(relu_op, {x}, tvm::Attrs(), {});
- auto z = relay::Call(add_op, {y, x}, tvm::Attrs(), {});
- auto ref_count = GetExprRefCount(z);
- ICHECK(ref_count[a.get()] == 1);
- ICHECK(ref_count[relu_op.get()] == 2);
- ICHECK(ref_count[add_op.get()] == 1);
- ICHECK(ref_count[x.get()] == 2);
- ICHECK(ref_count[y.get()] == 1);
- ICHECK(ref_count[z.get()] == 1);
-}
diff --git a/tests/cpp/relay_dismantler_test.cc b/tests/cpp/relay_dismantler_test.cc
deleted file mode 100644
index ae95185cb2..0000000000
--- a/tests/cpp/relay_dismantler_test.cc
+++ /dev/null
@@ -1,164 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-#include <gtest/gtest.h>
-#include <tvm/ir/expr.h>
-#include <tvm/ir/type_functor.h>
-#include <tvm/node/functor.h>
-#include <tvm/node/structural_equal.h>
-#include <tvm/relay/adt.h>
-#include <tvm/relay/analysis.h>
-#include <tvm/relay/expr.h>
-#include <tvm/relay/expr_functor.h>
-#include <tvm/relay/function.h>
-#include <tvm/relay/op.h>
-#include <tvm/relay/op_attr_types.h>
-#include <tvm/relay/op_strategy.h>
-#include <tvm/relay/transform.h>
-#include <tvm/relay/type.h>
-#include <tvm/runtime/packed_func.h>
-#include <tvm/runtime/registry.h>
-#include <tvm/te/operation.h>
-#include <tvm/topi/broadcast.h>
-#include <tvm/topi/generic/injective.h>
-
-#include <memory>
-
-using namespace tvm;
-using namespace tvm::relay;
-
-TEST(Relay, OutOfStack_add) {
- auto foo = [] {
- auto add_op = relay::Op::Get("add");
- auto c_data = tvm::runtime::NDArray::Empty({1, 2, 3}, {kDLFloat, 32, 1}, {kDLCPU, 0});
- auto c1 = relay::Constant(c_data);
- Call y1 = relay::Call(add_op, {c1, c1});
- for (int i = 0; i < 1e6; i++) {
- y1 = relay::Call(add_op, {c1, y1});
- }
- relay::Function func = relay::Function({}, y1, relay::Type(), {});
- };
- ASSERT_EXIT((foo(), exit(0)), ::testing::ExitedWithCode(0), ".*");
-}
-
-TEST(Relay, OutOfStack_cast) {
- auto foo = [] {
- auto cast_op = relay::Op::Get("cast");
- auto c_data = tvm::runtime::NDArray::Empty({1, 2, 3}, {kDLFloat, 32, 1}, {kDLCPU, 0});
- auto c1 = relay::Constant(c_data);
- Call y1 = relay::Call(cast_op, {c1});
- for (int i = 0; i < 1e6; i++) {
- y1 = relay::Call(cast_op, {y1});
- }
- relay::Function func = relay::Function({}, y1, relay::Type(), {});
- };
- ASSERT_EXIT((foo(), exit(0)), ::testing::ExitedWithCode(0), ".*");
-}
-
-TEST(Relay, OutOfStack_packed_func) {
- constexpr int len = 1e6;
- auto foo = [] {
- auto x = relay::Var("x", relay::TensorType({3, 2}, DataType::Float(32)));
- auto one = relay::Constant(tvm::runtime::NDArray::Empty({1}, {kDLFloat, 32, 1}, {kDLCPU, 0}));
- auto add_func = tvm::runtime::Registry::Get("relay.op._make.add");
- auto y = (*add_func)(x, one);
- for (int i = 0; i < len; ++i) {
- y = (*add_func)(y, one);
- }
-
- // check if still reachable
- int k = 0;
- Expr e = y;
- while (e.defined() && e.as<CallNode>() != nullptr) {
- e = e.as<CallNode>()->args[0];
- ++k;
- }
- ASSERT_EQ(len + 1, k);
- };
- ASSERT_EXIT((foo(), exit(0)), ::testing::ExitedWithCode(0), ".*");
-}
-
-TEST(Relay, CallNodeSharedArgs) {
- auto x = relay::Var("x", relay::TensorType({3, 2}, DataType::Float(32)));
- auto one = relay::Constant(tvm::runtime::NDArray::Empty({1}, {kDLFloat, 32, 1}, {kDLCPU, 0}));
- auto relu_op = relay::Op::Get("nn.relu");
- Call y = relay::Call(relu_op, {x}, Attrs(), {});
- y = relay::Call(relu_op, {y}, Attrs(), {});
- ASSERT_EQ(1, y.get()->args[0].as<CallNode>()->args.size());
- y = relay::Call(y.get()->op, y.get()->args, y.get()->attrs, y.get()->type_args);
- ASSERT_EQ(1, y.get()->args[0].as<CallNode>()->args.size());
-}
-
-TEST(Relay, TupleSharedFields) {
- auto x = relay::Var("x", relay::TensorType({3, 2}, DataType::Float(32)));
- auto one = relay::Constant(tvm::runtime::NDArray::Empty({1}, {kDLFloat, 32, 1}, {kDLCPU, 0}));
- auto relu_op = relay::Op::Get("nn.relu");
- Expr y = relay::Call(relu_op, {x}, Attrs(), {});
- y = relay::Call(relu_op, {y}, Attrs(), {});
- {
- Expr y1 = relay::Tuple(y.as<CallNode>()->args);
- Expr y2 = relay::Tuple(y.as<CallNode>()->args);
-
- y1 = relay::Call(relu_op, {y1});
- y2 = relay::Call(relu_op, {y2});
- y = y1;
- }
- ASSERT_EQ(1, y.as<CallNode>()->args[0].as<TupleNode>()->fields[0].as<CallNode>()->args.size());
-}
-
-TEST(Relay, TupleiGetItemSharedTuple) {
- auto x = relay::Var("x", relay::TensorType({3, 2}, DataType::Float(32)));
- auto one = relay::Constant(tvm::runtime::NDArray::Empty({1}, {kDLFloat, 32, 1}, {kDLCPU, 0}));
- auto relu_op = relay::Op::Get("nn.relu");
- Expr y = relay::Call(relu_op, {x}, Attrs(), {});
- y = relay::Tuple({y});
- {
- Expr y1 = relay::TupleGetItem(y, 0);
- Expr y2 = relay::TupleGetItem(y, 0);
-
- y1 = relay::Call(relu_op, {y1});
- y2 = relay::Call(relu_op, {y2});
- y = y1;
- }
- ASSERT_EQ(1, y.as<CallNode>()
- ->args[0]
- .as<TupleGetItemNode>()
- ->tuple.as<TupleNode>()
- ->fields[0]
- .as<CallNode>()
- ->args.size());
-}
-
-TEST(Relay, OutOfStackLet) {
- auto foo = [] {
- auto add_op = relay::Op::Get("add");
- auto p = relay::Var("p", relay::TensorType({3, 2}, DataType::Float(32)));
- int size = 1e6 - 1;
- std::vector<relay::Var> vars;
- for (int i = 0; i < size; ++i) {
- vars.emplace_back("x_" + std::to_string(i), relay::TensorType({3, 2}, DataType::Float(32)));
- }
- Expr body = vars[size - 1];
- for (int i = size - 1; i >= 0; --i) {
- Var v = i == 0 ? p : vars[i - 1];
- body = relay::Let(vars[i], relay::Call(add_op, {v, v}), body);
- }
- relay::Function func = relay::Function({p}, body, relay::Type(), {});
- };
- ASSERT_EXIT((foo(), exit(0)), ::testing::ExitedWithCode(0), ".*");
-}
diff --git a/tests/cpp/relay_pass_type_infer_test.cc b/tests/cpp/relay_pass_type_infer_test.cc
deleted file mode 100644
index 6db5952818..0000000000
--- a/tests/cpp/relay_pass_type_infer_test.cc
+++ /dev/null
@@ -1,43 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include <gtest/gtest.h>
-#include <tvm/node/structural_equal.h>
-#include <tvm/relay/analysis.h>
-#include <tvm/relay/expr.h>
-#include <tvm/relay/transform.h>
-#include <tvm/relay/type.h>
-#include <tvm/te/operation.h>
-
-TEST(Relay, SelfReference) {
- using namespace tvm;
- auto tensor_type = relay::TensorType({}, DataType::Bool());
- auto x = relay::Var("x", relay::Type());
- auto f = relay::Function(tvm::Array<relay::Var>{x}, x, relay::Type(), {});
- ICHECK(f->IsInstance<BaseFuncNode>());
- auto y = relay::Var("y", tensor_type);
- auto call = relay::Call(f, Array<relay::Expr>{y});
- auto fx = relay::Function(tvm::Array<relay::Var>{y}, call, relay::Type(), {});
- auto mod = IRModule::FromExpr(fx);
- mod = relay::transform::InferType()(mod);
- auto type_fx = mod->Lookup("main");
-
- auto expected = relay::FuncType(tvm::Array<relay::Type>{tensor_type}, tensor_type, {}, {});
- ICHECK(tvm::StructuralEqual()(type_fx->checked_type(), expected));
-}
diff --git a/tests/cpp/relay_text_printer_test.cc b/tests/cpp/relay_text_printer_test.cc
deleted file mode 100644
index 58fa228f8a..0000000000
--- a/tests/cpp/relay_text_printer_test.cc
+++ /dev/null
@@ -1,58 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include <gtest/gtest.h>
-#include <tvm/ir/expr.h>
-#include <tvm/ir/type_functor.h>
-#include <tvm/node/functor.h>
-#include <tvm/node/structural_equal.h>
-#include <tvm/relay/adt.h>
-#include <tvm/relay/analysis.h>
-#include <tvm/relay/expr.h>
-#include <tvm/relay/expr_functor.h>
-#include <tvm/relay/function.h>
-#include <tvm/relay/op.h>
-#include <tvm/relay/op_attr_types.h>
-#include <tvm/relay/op_strategy.h>
-#include <tvm/relay/transform.h>
-#include <tvm/relay/type.h>
-#include <tvm/runtime/packed_func.h>
-#include <tvm/runtime/registry.h>
-#include <tvm/te/operation.h>
-#include <tvm/topi/broadcast.h>
-#include <tvm/topi/generic/injective.h>
-
-using namespace tvm;
-using namespace tvm::relay;
-
-TEST(Relay, LargeGraphPrint) {
- auto foo = [] {
- auto add_op = relay::Op::Get("add");
- auto c_data = tvm::runtime::NDArray::Empty({1, 2, 3}, {kDLFloat, 32, 1}, {kDLCPU, 0});
- auto c1 = relay::Constant(c_data);
- Call y1 = relay::Call(add_op, {c1, c1});
- for (int i = 0; i < 1e6; i++) {
- y1 = relay::Call(add_op, {c1, y1});
- }
- relay::Function func = relay::Function({}, y1, relay::Type(), {});
- std::string result = AsText(func);
- ASSERT_GT(0, result.size());
- };
- ASSERT_EXIT((foo(), exit(0)), ::testing::ExitedWithCode(0), ".*");
-}
diff --git a/tests/cpp/relay_transform_sequential_test.cc b/tests/cpp/relay_transform_sequential_test.cc
deleted file mode 100644
index c4e5db89de..0000000000
--- a/tests/cpp/relay_transform_sequential_test.cc
+++ /dev/null
@@ -1,135 +0,0 @@
-/*
- * Licensed to the Apache Software Foundation (ASF) under one
- * or more contributor license agreements. See the NOTICE file
- * distributed with this work for additional information
- * regarding copyright ownership. The ASF licenses this file
- * to you under the Apache License, Version 2.0 (the
- * "License"); you may not use this file except in compliance
- * with the License. You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing,
- * software distributed under the License is distributed on an
- * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
- * KIND, either express or implied. See the License for the
- * specific language governing permissions and limitations
- * under the License.
- */
-
-#include <gtest/gtest.h>
-#include <tvm/driver/driver_api.h>
-#include <tvm/ir/module.h>
-#include <tvm/node/structural_equal.h>
-#include <tvm/relay/analysis.h>
-#include <tvm/relay/expr.h>
-#include <tvm/relay/op_attr_types.h>
-#include <tvm/relay/op_strategy.h>
-#include <tvm/relay/transform.h>
-#include <tvm/relay/type.h>
-#include <tvm/runtime/packed_func.h>
-#include <tvm/runtime/registry.h>
-#include <tvm/te/operation.h>
-#include <tvm/topi/broadcast.h>
-#include <tvm/topi/generic/injective.h>
-
-using namespace tvm;
-
-TVM_REGISTER_GLOBAL("test.seq.strategy")
- .set_body_typed([](const Attrs& attrs, const Array<te::Tensor>& inputs, const Type& out_type,
- const Target& target) {
- relay::FTVMCompute fcompute = [](const Attrs& attrs, const Array<te::Tensor>& inputs,
- const Type& out_type) -> Array<te::Tensor> {
- ICHECK_EQ(inputs.size(), 2U);
- return {topi::add(inputs[0], inputs[1])};
- };
- relay::FTVMSchedule fschedule = [](const Attrs& attrs, const Array<te::Tensor>& outs,
- const Target& target) {
- With<Target> target_scope(target);
- return topi::generic::schedule_injective(target, outs);
- };
-
- auto n = make_object<relay::OpStrategyNode>();
- auto strategy = relay::OpStrategy(std::move(n));
- strategy.AddImplementation(fcompute, fschedule, "test.strategy", 10);
- return strategy;
- });
-
-TEST(Relay, Sequential) {
- auto tensor_type = relay::TensorType({1, 2, 3}, DataType::Float(32));
- auto c_data = tvm::runtime::NDArray::Empty({1, 2, 3}, {kDLFloat, 32, 1}, {kDLCPU, 0});
-
- // Create a function for optimization.
- auto c = relay::Constant(c_data);
- auto a = relay::Var("a", tensor_type);
- auto x = relay::Var("x", tensor_type);
- auto add_op = relay::Op::Get("add");
- auto y = relay::Call(add_op, {c, c});
- y = relay::Call(add_op, {x, y});
- auto z = relay::Call(add_op, {y, c});
- auto z1 = relay::Call(add_op, {y, c});
- auto z2 = relay::Call(add_op, {z, z1});
- // Let expression and varaible a should be dead-code eliminated.
- auto z3 = relay::Let(a, c, z2);
- relay::Function func = relay::Function(relay::FreeVars(z3), z3, relay::Type(), {});
-
- auto reg = tvm::runtime::Registry::Get("ir.RegisterOpAttr");
- if (!reg) {
- LOG(FATAL) << "Register is not defined.";
- }
- auto reset = tvm::runtime::Registry::Get("ir.OpResetAttr");
- if (!reset) {
- LOG(FATAL) << "Reset is not defined.";
- }
- auto fs = tvm::runtime::Registry::Get("test.seq.strategy");
- if (!fs) {
- LOG(FATAL) << "Strategy is not defined.";
- }
- auto fgeneric = GenericFunc::Get("test.strategy_generic").set_default(*fs, true);
- (*reset)(add_op, "FTVMStrategy");
- (*reg)("add", "FTVMStrategy", fgeneric, 10);
-
- // Run sequential passes.
- tvm::Array<relay::transform::Pass> pass_seqs{
- relay::transform::InferType(), relay::transform::DeadCodeElimination(),
- relay::transform::EliminateCommonSubexpr(), relay::transform::AlterOpLayout()};
- relay::transform::Pass seq = relay::transform::Sequential(pass_seqs);
- auto mod = IRModule::FromExpr(func);
- auto pass_ctx = relay::transform::PassContext::Create();
- pass_ctx->opt_level = 3;
- pass_ctx->config.Set("relay.fallback_device_type", Integer(1));
- {
- tvm::With<relay::transform::PassContext> ctx_scope(pass_ctx);
- tvm::With<tvm::Target> tctx(tvm::Target("llvm"));
- mod = seq(mod);
- }
-
- ICHECK(mod.defined());
- auto entry_func = mod->GetGlobalVar("main");
- ICHECK(entry_func.defined());
- relay::Function f = Downcast<relay::Function>(mod->Lookup("main"));
- ICHECK(f.defined());
-
- // Expected function
- auto c1 = relay::Constant(c_data);
- auto x1 = relay::Var("x", tensor_type);
- auto y1 = relay::Call(add_op, {c1, c1});
- y1 = relay::Call(add_op, {x1, y1});
- auto zz = relay::Call(add_op, {y1, c1});
- zz = relay::Call(add_op, {zz, zz});
- relay::Function expected_func = relay::Function(relay::FreeVars(zz), zz, relay::Type(), {});
-
- // Infer type for the expected function.
- auto mod1 = IRModule::FromExpr(expected_func);
- mod1 = relay::transform::InferType()(mod1);
- auto expected = mod1->Lookup("main");
- ICHECK(tvm::StructuralEqual()(f, expected));
-}
-
-TEST(PassContextListConfigs, Basic) {
- Map<String, Map<String, String>> configs = relay::transform::PassContext::ListConfigs();
- ICHECK_EQ(configs.empty(), false);
-
- auto config = configs["relay.backend.use_auto_scheduler"];
- ICHECK_EQ(config["type"], "IntImm");
-}
diff --git a/tests/python/contrib/test_arm_compute_lib/__init__.py b/tests/python/contrib/test_arm_compute_lib/__init__.py
deleted file mode 100644
index fd14be1cc3..0000000000
--- a/tests/python/contrib/test_arm_compute_lib/__init__.py
+++ /dev/null
@@ -1,17 +0,0 @@
-# Licensed to the Apache Software Foundation (ASF) under one
-# or more contributor license agreements. See the NOTICE file
-# distributed with this work for additional information
-# regarding copyright ownership. The ASF licenses this file
-# to you under the Apache License, Version 2.0 (the
-# "License"); you may not use this file except in compliance
-# with the License. You may obtain a copy of the License at
-#
-# http://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing,
-# software distributed under the License is distributed on an
-# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
-# KIND, either express or implied. See the License for the
-# specific language governing permissions and limitations
-# under the License.
-"""Infrastructure and tests for Arm Compute Library"""
diff --git a/tests/python/contrib/test_arm_compute_lib/infrastructure.py b/tests/python/contrib/test_arm_compute_lib/infrastructure.py
deleted file mode 100644
index ae3acbb096..0000000000
--- a/tests/python/contrib/test_arm_compute_lib/infrastructure.py
+++ /dev/null
@@ -1,341 +0,0 @@
-# Licensed to the Apache Software Foundation (ASF) under one
-# or more contributor license agreements. See the NOTICE file
-# distributed with this work for additional information
-# regarding copyright ownership. The ASF licenses this file
-# to you under the Apache License, Version 2.0 (the
-# "License"); you may not use this file except in compliance
-# with the License. You may obtain a copy of the License at
-#
-# http://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing,
-# software distributed under the License is distributed on an
-# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
-# KIND, either express or implied. See the License for the
-# specific language governing permissions and limitations
-# under the License.
-from itertools import zip_longest, combinations
-import json
-import os
-import warnings
-
-import numpy as np
-
-import tvm
-from tvm import relay
-from tvm import rpc
-from tvm.contrib import graph_executor
-from tvm.relay.op.contrib import arm_compute_lib
-from tvm.contrib import utils
-from tvm.autotvm.measure import request_remote
-
-
-QNN_DTYPES = ("uint8", "int8")
-
-
-class Device:
- """
- Configuration for Arm Compute Library tests.
-
- Check tests/python/contrib/arm_compute_lib/ for the presence of an test_config.json file.
- This file can be used to override the default configuration here which will attempt to run the Arm
- Compute Library runtime tests locally if the runtime is available. Changing the configuration
- will allow these runtime tests to be offloaded to a remote Arm device via a tracker for example.
-
- Notes
- -----
- The test configuration will be loaded once when the class is created. If the configuration
- changes between tests, any changes will not be picked up.
-
- Parameters
- ----------
- device : RPCSession
- Allows tests to connect to and use remote device.
-
- Attributes
- ----------
- connection_type : str
- Details the type of RPC connection to use. Options:
- local - Use the local device,
- tracker - Connect to a tracker to request a remote device,
- remote - Connect to a remote device directly.
- host : str
- Specify IP address or hostname of remote target.
- port : int
- Specify port number of remote target.
- target : str
- The compilation target.
- device_key : str
- The device key of the remote target. Use when connecting to a remote device via a tracker.
- cross_compile : str
- Specify path to cross compiler to use when connecting a remote device from a non-arm platform.
- """
-
- connection_type = "local"
- host = "127.0.0.1"
- port = 9090
- target = "llvm -mtriple=aarch64-linux-gnu -mattr=+neon"
- device_key = ""
- cross_compile = ""
-
- def __init__(self):
- """Keep remote device for lifetime of object."""
- self.device = self._get_remote()
-
- @classmethod
- def _get_remote(cls):
- """Get a remote (or local) device to use for testing."""
- if cls.connection_type == "tracker":
- device = request_remote(cls.device_key, cls.host, cls.port, timeout=1000)
- elif cls.connection_type == "remote":
- device = rpc.connect(cls.host, cls.port)
- elif cls.connection_type == "local":
- device = rpc.LocalSession()
- else:
- raise ValueError(
- "connection_type in test_config.json should be one of: " "local, tracker, remote."
- )
-
- return device
-
- @classmethod
- def load(cls, file_name):
- """Load test config
-
- Load the test configuration by looking for file_name relative
- to the test_arm_compute_lib directory.
- """
- location = os.path.realpath(os.path.join(os.getcwd(), os.path.dirname(__file__)))
- config_file = os.path.join(location, file_name)
- if not os.path.exists(config_file):
- warnings.warn(
- "Config file doesn't exist, resuming Arm Compute Library tests with default config."
- )
- return
- with open(config_file, mode="r") as config:
- test_config = json.load(config)
-
- cls.connection_type = test_config["connection_type"]
- cls.host = test_config["host"]
- cls.port = test_config["port"]
- cls.target = test_config["target"]
- cls.device_key = test_config.get("device_key") or ""
- cls.cross_compile = test_config.get("cross_compile") or ""
-
-
-def get_low_high_atol_rtol(dtype):
- """Returns a tuple with boundary values and tolerance for ACL tests."""
-
- if dtype == "float32":
- low, high, atol, rtol = (-127, 128, 0.001, 0.001)
- elif dtype == "uint8":
- low, high, atol, rtol = (0, 255, 1, 0)
- elif dtype == "int8":
- low, high, atol, rtol = (-127, 128, 1, 0)
- else:
- raise Exception(f"dtype not expected: {dtype}")
-
- return low, high, atol, rtol
-
-
-def get_cpu_op_count(mod):
- """Traverse graph counting ops offloaded to TVM."""
-
- class Counter(tvm.relay.ExprVisitor):
- def __init__(self):
- super().__init__()
- self.count = 0
-
- def visit_call(self, call):
- if isinstance(call.op, tvm.ir.Op):
- self.count += 1
-
- super().visit_call(call)
-
- c = Counter()
- c.visit(mod["main"])
- return c.count
-
-
-def skip_runtime_test():
- """Skip test if it requires the runtime and it's not present."""
- # ACL codegen not present.
- if not tvm.get_global_func("relay.ext.arm_compute_lib", True):
- print("Skip because Arm Compute Library codegen is not available.")
- return True
-
- # Remote device is in use or ACL runtime not present
- # Note: Ensure that the device config has been loaded before this check
- if (
- not Device.connection_type != "local"
- and not arm_compute_lib.is_arm_compute_runtime_enabled()
- ):
- print("Skip because runtime isn't present or a remote device isn't being used.")
- return True
-
-
-def skip_codegen_test():
- """Skip test if it requires the ACL codegen and it's not present."""
- if not tvm.get_global_func("relay.ext.arm_compute_lib", True):
- print("Skip because Arm Compute Library codegen is not available.")
- return True
-
-
-def build_module(
- mod,
- target,
- params=None,
- enable_acl=True,
- tvm_ops=0,
- acl_partitions=1,
- disabled_ops=["concatenate"],
-):
- """Build module with option to build for ACL."""
- if isinstance(mod, tvm.relay.expr.Call):
- mod = tvm.IRModule.from_expr(mod)
- with tvm.transform.PassContext(opt_level=3, disabled_pass=["AlterOpLayout"]):
- if enable_acl:
- mod = arm_compute_lib.partition_for_arm_compute_lib(
- mod, params, disabled_ops=disabled_ops
- )
- tvm_op_count = get_cpu_op_count(mod)
- assert tvm_op_count == tvm_ops, "Got {} TVM operators, expected {}".format(
- tvm_op_count, tvm_ops
- )
- partition_count = 0
- for global_var in mod.get_global_vars():
- if "arm_compute_lib" in global_var.name_hint:
- partition_count += 1
-
- assert (
- acl_partitions == partition_count
- ), "Got {} Arm Compute Library partitions, expected {}".format(
- partition_count, acl_partitions
- )
- relay.backend.te_compiler.get().clear()
- return relay.build(mod, target=target, params=params)
-
-
-def build_and_run(
- mod,
- inputs,
- outputs,
- params,
- device,
- enable_acl=True,
- no_runs=1,
- tvm_ops=0,
- acl_partitions=1,
- config=None,
- disabled_ops=["concatenate"],
-):
- """Build and run the relay module."""
- if config is None:
- config = {}
-
- try:
- lib = build_module(
- mod, device.target, params, enable_acl, tvm_ops, acl_partitions, disabled_ops
- )
- except Exception as e:
- err_msg = "The module could not be built.\n"
- if config:
- err_msg += f"The test failed with the following parameters: {config}\n"
- err_msg += str(e)
- raise Exception(err_msg)
-
- lib = update_lib(lib, device.device, device.cross_compile)
- gen_module = graph_executor.GraphModule(lib["default"](device.device.cpu(0)))
- gen_module.set_input(**inputs)
- out = []
- for _ in range(no_runs):
- gen_module.run()
- out.append([gen_module.get_output(i) for i in range(outputs)])
- return out
-
-
-def update_lib(lib, device, cross_compile):
- """Export the library to the remote/local device."""
- lib_name = "mod.so"
- temp = utils.tempdir()
- lib_path = temp.relpath(lib_name)
- if cross_compile:
- lib.export_library(lib_path, cc=cross_compile)
- else:
- lib.export_library(lib_path)
- device.upload(lib_path)
- lib = device.load_module(lib_name)
- return lib
-
-
-def verify(answers, atol, rtol, verify_saturation=False, config=None):
- """Compare the array of answers. Each entry is a list of outputs."""
- if config is None:
- config = {}
-
- if len(answers) < 2:
- raise RuntimeError(f"No results to compare: expected at least two, found {len(answers)}")
- for answer in zip_longest(*answers):
- for outs in combinations(answer, 2):
- try:
- if verify_saturation:
- assert (
- np.count_nonzero(outs[0].numpy() == 255) < 0.25 * outs[0].numpy().size
- ), "Output is saturated: {}".format(outs[0])
- assert (
- np.count_nonzero(outs[0].numpy() == 0) < 0.25 * outs[0].numpy().size
- ), "Output is saturated: {}".format(outs[0])
- tvm.testing.assert_allclose(outs[0].numpy(), outs[1].numpy(), rtol=rtol, atol=atol)
- except AssertionError as e:
- err_msg = "Results not within the acceptable tolerance.\n"
- if config:
- err_msg += f"The test failed with the following parameters: {config}\n"
- err_msg += str(e)
- raise AssertionError(err_msg)
-
-
-def extract_acl_modules(module):
- """Get the ACL module(s) from llvm module."""
- return list(
- filter(lambda mod: mod.type_key == "arm_compute_lib", module.get_lib().imported_modules)
- )
-
-
-def verify_codegen(
- module,
- known_good_codegen,
- num_acl_modules=1,
- tvm_ops=0,
- target="llvm -mtriple=aarch64-linux-gnu -mattr=+neon",
- disabled_ops=["concatenate"],
-):
- """Check acl codegen against a known good output."""
- module = build_module(
- module,
- target,
- tvm_ops=tvm_ops,
- acl_partitions=num_acl_modules,
- disabled_ops=disabled_ops,
- )
- acl_modules = extract_acl_modules(module)
-
- assert len(acl_modules) == num_acl_modules, (
- f"The number of Arm Compute Library modules produced ({len(acl_modules)}) does not "
- f"match the expected value ({num_acl_modules})."
- )
-
- for mod in acl_modules:
- source = mod.get_source("json")
- codegen = json.loads(source)["nodes"]
- # remove input and const names as these cannot be predetermined
- for node in range(len(codegen)):
- if codegen[node]["op"] == "input" or codegen[node]["op"] == "const":
- codegen[node]["name"] = ""
- codegen_str = json.dumps(codegen, sort_keys=True, indent=2)
- known_good_codegen_str = json.dumps(known_good_codegen, sort_keys=True, indent=2)
-
- assert codegen_str == known_good_codegen_str, (
- f"The JSON produced by codegen does not match the expected result. \n"
- f"Actual={codegen_str} \n"
- f"Expected={known_good_codegen_str}"
- )
diff --git a/tests/python/contrib/test_arm_compute_lib/test_add.py b/tests/python/contrib/test_arm_compute_lib/test_add.py
deleted file mode 100644
index 319105bb5f..0000000000
--- a/tests/python/contrib/test_arm_compute_lib/test_add.py
+++ /dev/null
@@ -1,168 +0,0 @@
-# Licensed to the Apache Software Foundation (ASF) under one
-# or more contributor license agreements. See the NOTICE file
-# distributed with this work for additional information
-# regarding copyright ownership. The ASF licenses this file
-# to you under the Apache License, Version 2.0 (the
-# "License"); you may not use this file except in compliance
-# with the License. You may obtain a copy of the License at
-#
-# http://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing,
-# software distributed under the License is distributed on an
-# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
-# KIND, either express or implied. See the License for the
-# specific language governing permissions and limitations
-# under the License.
-"""Arm Compute Library integration reshape tests."""
-
-import numpy as np
-import pytest
-
-import tvm
-import tvm.testing
-from tvm import relay
-
-from test_arm_compute_lib.infrastructure import (
- skip_runtime_test,
- skip_codegen_test,
- build_and_run,
- verify,
- verify_codegen,
-)
-from test_arm_compute_lib.infrastructure import Device
-
-_qnn_params = {
- "lhs_scale": relay.const(0.0156863, "float32"),
- "lhs_zero_point": relay.const(127, "int32"),
- "rhs_scale": relay.const(0.0117647, "float32"),
- "rhs_zero_point": relay.const(85, "int32"),
- "output_scale": relay.const(0.0235294, "float32"),
- "output_zero_point": relay.const(128, "int32"),
-}
-
-
-def _get_model(shape, dtype, var_names, op, op_params):
- a = relay.var(next(var_names), shape=shape, dtype=dtype)
- b = relay.var(next(var_names), shape=shape, dtype=dtype)
- return op(a, b, **op_params)
-
-
-def _get_expected_codegen(shape, dtype, op_name, qnn_params):
- input_a = {"op": "input", "name": "", "attrs": {"shape": [[list(shape)]], "dtype": [[dtype]]}}
- input_b = {"op": "input", "name": "", "attrs": {"shape": [[list(shape)]], "dtype": [[dtype]]}}
- input_qnn = [
- {
- "op": "const",
- "name": "",
- "attrs": {
- "shape": [[list(qnn_params[_].data.shape)]],
- "dtype": [[qnn_params[_].data.dtype]],
- },
- }
- for _ in qnn_params
- ]
- inputs = [input_a, input_b, *input_qnn]
- node = {
- "op": "kernel",
- "name": op_name,
- "inputs": [[_, 0, 0] for _ in range(len(inputs))],
- "attrs": {
- "num_inputs": str(len(inputs)),
- "num_outputs": "1",
- "shape": [[list(shape)]],
- "dtype": [[dtype]],
- },
- }
-
- if qnn_params:
- node["attrs"]["lhs_axis"] = [["-1"]]
- node["attrs"]["rhs_axis"] = [["-1"]]
-
- return [*inputs, node]
-
-
-def test_runtime_add():
- Device.load("test_config.json")
-
- if skip_runtime_test():
- return
-
- device = Device()
- np.random.seed(0)
-
- for dtype, low, high, atol, rtol, op, op_params in [
- ("float32", -127, 128, 1e-7, 1e-7, relay.add, {}),
- ("uint8", 0, 255, 1.0, 0.0, relay.qnn.op.add, _qnn_params),
- ("int8", -127, 128, 1.0, 0.0, relay.qnn.op.add, _qnn_params),
- ]:
- shape = (2, 2)
- for inputs in [
- {
- "a": tvm.nd.array(np.random.uniform(low, high, shape).astype(dtype)),
- "b": tvm.nd.array(np.random.uniform(low, high, shape).astype(dtype)),
- }
- ]:
- outputs = []
- func = _get_model(shape, dtype, iter(inputs), op, op_params)
- for acl in [True, False]:
- outputs.append(build_and_run(func, inputs, 1, None, device, enable_acl=acl)[0])
-
- config = {
- "shape": shape,
- "dtype": dtype,
- "inputs": inputs,
- "operation": op,
- "op_params": op_params,
- }
-
- verify(outputs, atol=atol, rtol=rtol, config=config, verify_saturation=False)
-
-
-def test_codegen_add():
- if skip_codegen_test():
- return
-
- inputs = {"a", "b"}
- for dtype, op_name, op, qnn_params in [
- ("float32", "add", relay.add, {}),
- ("uint8", "qnn.add", relay.qnn.op.add, _qnn_params),
- ("int8", "qnn.add", relay.qnn.op.add, _qnn_params),
- ]:
- for shape in [(1, 1), (2, 2, 2), (3, 3, 3, 3)]:
- func = _get_model(shape, dtype, iter(inputs), op, qnn_params)
- exp_codegen = _get_expected_codegen(shape, dtype, op_name, qnn_params)
- verify_codegen(func, exp_codegen, 1)
-
-
-@pytest.mark.parametrize(
- "param, param_type",
- [
- ("lhs_scale", "float32"),
- ("lhs_zero_point", "int32"),
- ("rhs_scale", "float32"),
- ("rhs_zero_point", "int32"),
- ],
-)
-def test_codegen_add_per_channel_quantization(param, param_type):
- if skip_codegen_test():
- return
-
- qnn_params = _qnn_params
- qnn_params[param] = relay.const([1, 2], param_type)
-
- dtype = "int8"
- op_name = "qnn.add"
- op = relay.qnn.op.add
- inputs = {"a", "b"}
-
- for shape in [(1, 3, 3, 2)]:
- func = _get_model(shape, dtype, iter(inputs), op, qnn_params)
- exp_codegen = _get_expected_codegen(shape, dtype, op_name, qnn_params)
- verify_codegen(func, exp_codegen, num_acl_modules=0, tvm_ops=1)
-
-
-if __name__ == "__main__":
- test_runtime_add()
- test_codegen_add()
- test_codegen_add_per_channel_quantization()
diff --git a/tests/python/contrib/test_arm_compute_lib/test_concatenate.py b/tests/python/contrib/test_arm_compute_lib/test_concatenate.py
deleted file mode 100644
index 55072f37c2..0000000000
--- a/tests/python/contrib/test_arm_compute_lib/test_concatenate.py
+++ /dev/null
@@ -1,159 +0,0 @@
-# Licensed to the Apache Software Foundation (ASF) under one
-# or more contributor license agreements. See the NOTICE file
-# distributed with this work for additional information
-# regarding copyright ownership. The ASF licenses this file
-# to you under the Apache License, Version 2.0 (the
-# "License"); you may not use this file except in compliance
-# with the License. You may obtain a copy of the License at
-#
-# http://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing,
-# software distributed under the License is distributed on an
-# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
-# KIND, either express or implied. See the License for the
-# specific language governing permissions and limitations
-# under the License.
-"""Arm Compute Library integration concatenate tests."""
-
-import numpy as np
-import pytest
-
-import tvm
-from tvm import relay
-from tvm import testing
-
-from test_arm_compute_lib.infrastructure import (
- skip_runtime_test,
- skip_codegen_test,
- build_and_run,
- verify,
- verify_codegen,
-)
-from test_arm_compute_lib.infrastructure import Device
-
-
-def _get_model(input_shape_a, input_shape_b, input_shape_c, axis, dtype, var_names):
- """Return a model and any parameters it may have."""
- a = relay.var(next(var_names), shape=input_shape_a, dtype=dtype)
- b = relay.var(next(var_names), shape=input_shape_b, dtype=dtype)
- c = relay.var(next(var_names), shape=input_shape_c, dtype=dtype)
- out = relay.concatenate([a, b, c], axis)
- return out
-
-
-def _get_expected_codegen(input_shape_a, input_shape_b, input_shape_c, axis, dtype):
- node = {
- "op": "kernel",
- "name": "concatenate",
- "inputs": [
- [0, 0, 0],
- [1, 0, 0],
- [2, 0, 0],
- ],
- "attrs": {
- "num_outputs": "1",
- "num_inputs": "3",
- "dtype": [[dtype]],
- "axis": [[str(axis)]],
- "shape": [[[6, 234, 234, 256]]],
- },
- }
-
- input_a = {
- "op": "input",
- "name": "",
- "attrs": {
- "shape": [[input_shape_a]],
- "dtype": [[dtype]],
- },
- }
-
- input_b = {
- "op": "input",
- "name": "",
- "attrs": {
- "shape": [[input_shape_b]],
- "dtype": [[dtype]],
- },
- }
-
- input_c = {
- "op": "input",
- "name": "",
- "attrs": {
- "shape": [[input_shape_c]],
- "dtype": [[dtype]],
- },
- }
- return [input_a, input_b, input_c, node]
-
-
-@pytest.mark.parametrize(
- "input_shape_a, input_shape_b, input_shape_c, axis, dtype",
- [
- ([1, 234, 234, 256], [2, 234, 234, 256], [3, 234, 234, 256], 0, "float32"),
- ([1, 1, 234, 256], [1, 2, 234, 256], [1, 3, 234, 256], 1, "float32"),
- ([1, 234, 234, 1], [1, 234, 234, 2], [1, 234, 234, 3], -1, "float32"),
- ([1, 234, 234, 256], [2, 234, 234, 256], [3, 234, 234, 256], -4, "float32"),
- ([1, 234, 234, 256], [2, 234, 234, 256], [3, 234, 234, 256], 0, "uint8"),
- ([1, 1, 234, 256], [1, 2, 234, 256], [1, 3, 234, 256], 1, "uint8"),
- ([1, 234, 234, 1], [1, 234, 234, 2], [1, 234, 234, 3], -1, "uint8"),
- ([1, 234, 234, 256], [2, 234, 234, 256], [3, 234, 234, 256], -4, "uint8"),
- ([1, 234, 234, 256], [2, 234, 234, 256], [3, 234, 234, 256], 0, "int8"),
- ([1, 1, 234, 256], [1, 2, 234, 256], [1, 3, 234, 256], 1, "int8"),
- ([1, 234, 234, 1], [1, 234, 234, 2], [1, 234, 234, 3], -1, "int8"),
- ([1, 234, 234, 256], [2, 234, 234, 256], [3, 234, 234, 256], -4, "int8"),
- ],
-)
-def test_concatenate(input_shape_a, input_shape_b, input_shape_c, axis, dtype):
- Device.load("test_config.json")
-
- if skip_runtime_test():
- return
-
- device = Device()
- np.random.seed(0)
-
- outputs = []
- inputs = {
- "a": tvm.nd.array(np.random.randn(*input_shape_a).astype(dtype)),
- "b": tvm.nd.array(np.random.randn(*input_shape_b).astype(dtype)),
- "c": tvm.nd.array(np.random.randn(*input_shape_c).astype(dtype)),
- }
- func = _get_model(
- inputs["a"].shape, inputs["b"].shape, inputs["c"].shape, axis, dtype, iter(inputs)
- )
- for acl in [False, True]:
- outputs.append(
- build_and_run(func, inputs, 1, None, device, enable_acl=acl, disabled_ops=[])[0]
- )
-
- config = {
- "input_shape_a": input_shape_a,
- "input_shape_b": input_shape_b,
- "input_shape_c": input_shape_c,
- "axis": axis,
- "dtype": dtype,
- }
- verify(outputs, atol=1e-7, rtol=1e-7, config=config)
-
-
-def test_codegen_concatenate():
- if skip_codegen_test():
- return
- shape_a = [1, 234, 234, 256]
- shape_b = [2, 234, 234, 256]
- shape_c = [3, 234, 234, 256]
- axis = 0
- inputs = {"a", "b", "c"}
- for dtype in ["float32"]:
- args = (shape_a, shape_b, shape_c, axis, dtype)
- func = _get_model(*args, iter(inputs))
- exp_codegen = _get_expected_codegen(*args)
- verify_codegen(func, exp_codegen, 1, disabled_ops=[])
-
-
-if __name__ == "__main__":
- test_concatenate()
- test_codegen_concatenate()
diff --git a/tests/python/contrib/test_arm_compute_lib/test_config.json b/tests/python/contrib/test_arm_compute_lib/test_config.json
deleted file mode 100644
index 5c75e659af..0000000000
--- a/tests/python/contrib/test_arm_compute_lib/test_config.json
+++ /dev/null
@@ -1,8 +0,0 @@
-{
- "connection_type": "local",
- "host": "127.0.0.1",
- "port": 9090,
- "target": "llvm -mtriple=aarch64-linux-gnu -mattr=+neon",
- "device_key": "",
- "cross_compile": ""
-}
diff --git a/tests/python/contrib/test_arm_compute_lib/test_conv2d.py b/tests/python/contrib/test_arm_compute_lib/test_conv2d.py
deleted file mode 100644
index b4fa49ffa2..0000000000
--- a/tests/python/contrib/test_arm_compute_lib/test_conv2d.py
+++ /dev/null
@@ -1,672 +0,0 @@
-# Licensed to the Apache Software Foundation (ASF) under one
-# or more contributor license agreements. See the NOTICE file
-# distributed with this work for additional information
-# regarding copyright ownership. The ASF licenses this file
-# to you under the Apache License, Version 2.0 (the
-# "License"); you may not use this file except in compliance
-# with the License. You may obtain a copy of the License at
-#
-# http://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing,
-# software distributed under the License is distributed on an
-# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
-# KIND, either express or implied. See the License for the
-# specific language governing permissions and limitations
-# under the License.
-"""Arm Compute Library integration conv2d tests."""
-
-import numpy as np
-import pytest
-
-import tvm
-from tvm import relay
-
-from test_arm_compute_lib.infrastructure import (
- QNN_DTYPES,
- get_low_high_atol_rtol,
- skip_runtime_test,
- skip_codegen_test,
- build_and_run,
- verify,
- verify_codegen,
-)
-from test_arm_compute_lib.infrastructure import Device
-
-
-def _get_model(
- shape,
- kernel_h,
- kernel_w,
- padding,
- strides,
- dilation,
- groups,
- dtype,
- channels,
- var_names,
- has_bias=False,
- has_activation=False,
- has_pad=False,
-):
- """Return a model and any parameters it may have"""
- a = relay.var(next(var_names), shape=shape, dtype=dtype)
- if has_pad:
- p = ((0, 0), (padding[0], padding[0]), (padding[1], padding[1]), (0, 0))
- a = relay.nn.pad(a, pad_width=p)
- padding = (0, 0, 0, 0)
- else:
- if len(padding) == 2:
- padding = (padding[0], padding[1], padding[0], padding[1])
- shape = (shape[0], shape[1] + padding[0] * 2, shape[2] + padding[1] * 2, shape[3])
- is_depthwise = shape[3] == channels == groups
- weight_format = "HWOI" if is_depthwise else "HWIO"
- if weight_format == "HWIO":
- weight_shape = (kernel_h, kernel_w, shape[3] // groups, channels)
- else:
- weight_shape = (kernel_h, kernel_w, channels, shape[3] // groups)
- w = tvm.nd.array(np.random.uniform(-128, 127, weight_shape).astype(dtype))
- weights = relay.const(w, dtype)
- out = relay.nn.conv2d(
- a,
- weights,
- kernel_size=(kernel_h, kernel_w),
- data_layout="NHWC",
- kernel_layout=weight_format,
- dilation=dilation,
- strides=strides,
- padding=padding,
- groups=groups,
- channels=channels,
- out_dtype=dtype,
- )
- params = {"w": w}
- if has_bias:
- bias_shape = weight_shape[2] if is_depthwise else weight_shape[3]
- b = tvm.nd.array(np.random.uniform(-128, 127, bias_shape).astype(dtype))
- biasc = relay.const(b, dtype)
- out = relay.nn.bias_add(out, biasc, axis=3)
- params["b"] = b
- if has_activation:
- out = relay.nn.relu(out)
- return out, params
-
-
-def _get_qnn_params(input_zp, input_sc, kernel_zp, kernel_sc, kernel_h, kernel_w, channels):
- """Get output qnn parameters given input and kernel parameters."""
- input_max = input_sc * (255 - input_zp)
- input_min = -input_sc * input_zp
- kernel_max = kernel_sc * (255 - kernel_zp)
- kernel_min = -kernel_sc * kernel_zp
- output_limits = [
- kernel_max * kernel_h * kernel_w * channels * input_max,
- kernel_min * kernel_h * kernel_w * channels * input_max,
- kernel_min * kernel_h * kernel_w * channels * input_min,
- kernel_max * kernel_h * kernel_w * channels * input_min,
- ]
- output_max = max(output_limits)
- output_min = min(output_limits)
- output_sc = (output_max - output_min) / 255
- output_zp = -int(output_min / output_sc)
- return output_zp, output_sc
-
-
-def _get_qnn_model(
- shape,
- kernel_h,
- kernel_w,
- padding,
- strides,
- dilation,
- groups,
- dtype,
- channels,
- input_zp,
- input_sc,
- kernel_zp,
- kernel_sc,
- output_zp,
- output_sc,
- var_names,
- has_bias=False,
- has_activation=False,
- has_pad=False,
-):
- """Return a model and any parameters it may have."""
- low, high, _, _ = get_low_high_atol_rtol(dtype)
-
- a = relay.var(next(var_names), shape=shape, dtype=dtype)
- if has_pad:
- p = ((0, 0), (padding[0], padding[0]), (padding[1], padding[1]), (0, 0))
- a = relay.nn.pad(a, pad_width=p, pad_value=input_zp, pad_mode="constant")
- padding = (0, 0, 0, 0)
- else:
- if len(padding) == 2:
- padding = (padding[0], padding[1], padding[0], padding[1])
- shape = (shape[0], shape[1] + padding[0] * 2, shape[2] + padding[1] * 2, shape[3])
- is_depthwise = shape[3] == channels == groups
- weight_format = "HWOI" if is_depthwise else "HWIO"
- if weight_format == "HWIO":
- weight_shape = (kernel_h, kernel_w, shape[3] // groups, channels)
- else:
- weight_shape = (kernel_h, kernel_w, channels, shape[3] // groups)
- w = tvm.nd.array(np.random.uniform(low, high, weight_shape).astype(dtype))
- weights = relay.const(w, dtype)
- out = relay.qnn.op.conv2d(
- a,
- weights,
- input_zero_point=relay.const(input_zp, "int32"),
- kernel_zero_point=relay.const(kernel_zp, "int32"),
- input_scale=relay.const(input_sc, "float32"),
- kernel_scale=relay.const(kernel_sc, "float32"),
- kernel_size=(kernel_h, kernel_w),
- data_layout="NHWC",
- kernel_layout=weight_format,
- dilation=dilation,
- strides=strides,
- padding=padding,
- groups=groups,
- channels=channels,
- out_dtype="int32",
- )
- params = {"w": w}
- if has_bias:
- bias_shape = weight_shape[2] if is_depthwise else weight_shape[3]
- b = tvm.nd.array(np.random.uniform(-128, 127, bias_shape).astype("int32"))
- biasc = relay.const(b, "int32")
- out = relay.nn.bias_add(out, biasc, axis=3)
- params["b"] = b
- if has_activation:
- out = relay.nn.relu(out)
- req = relay.qnn.op.requantize(
- out,
- relay.const(input_sc * kernel_sc, "float32"), # input scale
- relay.const(0, "int32"), # input zero point
- relay.const(output_sc, "float32"), # output scale
- relay.const(output_zp, "int32"), # output zero point
- out_dtype=dtype,
- )
- return req, params
-
-
-def _get_expected_codegen(
- shape,
- kernel_h,
- kernel_w,
- padding,
- strides,
- dilation,
- groups,
- dtype,
- channels,
- has_bias=False,
- has_activation=False,
-):
- if len(padding) == 2:
- padding = (padding[0], padding[1], padding[0], padding[1])
- output_height = ((shape[1] - kernel_h + padding[0] + padding[2]) / strides[0]) + 1
- output_width = ((shape[2] - kernel_w + padding[1] + padding[3]) / strides[1]) + 1
- output_shape = (1, int(output_height), int(output_width), channels)
- out_dtype = "int32" if dtype in QNN_DTYPES else "float32"
- is_depthwise = shape[3] == channels == groups
- weight_format = "IHWO" if is_depthwise else "OHWI"
- if weight_format == "IHWO":
- weight_shape = (shape[3] // groups, kernel_h, kernel_w, channels)
- else:
- weight_shape = (channels, kernel_h, kernel_w, shape[3] // groups)
- if is_depthwise:
- name = "nn.depthwise_conv2d"
- else:
- name = "nn.conv2d"
-
- node = {
- "op": "kernel",
- "name": name,
- "inputs": [],
- "attrs": {
- "groups": [[str(groups)]],
- "num_outputs": "1",
- "data_layout": [["NHWC"]],
- "kernel_layout": [[weight_format]],
- "channels": [[str(channels)]],
- "dilation": [[str(dilation[0]), str(dilation[1])]],
- "out_layout": [[""]],
- "out_dtype": [[out_dtype]],
- "kernel_size": [[str(kernel_h), str(kernel_w)]],
- "shape": [[list(output_shape)]],
- "dtype": [[dtype]],
- "padding": [[str(p) for p in padding]],
- "strides": [[str(s) for s in strides]],
- },
- }
-
- if has_activation:
- node["attrs"]["activation_type"] = [["relu"]]
-
- inputs = [
- {"op": "input", "name": "", "attrs": {"shape": [[list(shape)]], "dtype": [[str(dtype)]]}},
- {
- "op": "const",
- "name": "",
- "attrs": {"shape": [[list(weight_shape)]], "dtype": [[str(dtype)]]},
- },
- ]
-
- # qnn.conv2d params, input and kernel
- if dtype in QNN_DTYPES:
- node["name"] = "qnn." + node["name"].split(".")[1]
- for param_dtype in ["int32", "float32"]:
- for _ in range(2):
- inputs.append(
- {
- "op": "const",
- "name": "",
- "attrs": {"shape": [[[]]], "dtype": [[param_dtype]]},
- }
- )
-
- if has_bias:
- bias_dtype = "int32" if dtype in QNN_DTYPES else "float32"
- inputs.append(
- {
- "op": "const",
- "name": "",
- "attrs": {
- "shape": [[[1, 1, 1, weight_shape[3] if is_depthwise else weight_shape[0]]]],
- "dtype": [[bias_dtype]],
- },
- }
- )
-
- # qnn.conv2d params, output
- if dtype in QNN_DTYPES:
- for param_dtype in ["float32", "int32"]:
- inputs.append(
- {"op": "const", "name": "", "attrs": {"shape": [[[]]], "dtype": [[param_dtype]]}}
- )
-
- input_idx = 0
- for _ in range(len(inputs)):
- node["inputs"].append([input_idx, 0, 0])
- input_idx += 1
- node["attrs"]["num_inputs"] = str(len(inputs))
- inputs.append(node)
- return inputs
-
-
-@pytest.mark.parametrize(
- "trial",
- [
- # Normal convolution
- [2, 2, (1, 1), (1, 1), (1, 1), 4, (10, 10, 14), (False, False, False), False],
- [2, 1, (2, 2), (1, 1), (1, 1), 7, (12, 15, 16), (False, False, True), False],
- [3, 3, (2, 1), (1, 1), (1, 1), 4, (10, 10, 14), (False, True, False), False],
- [3, 3, (1, 1), (1, 1), (1, 1), 16, (12, 15, 16), (False, False, False), False],
- [5, 5, (1, 1), (2, 2), (1, 1), 4, (10, 10, 14), (True, False, False), False],
- [1, 3, (1, 1), (1, 1), (1, 1), 7, (20, 20, 20), (False, False, True), False],
- [2, 2, (2, 2), (1, 1), (1, 1), 4, (20, 20, 20), (False, True, False), False],
- [5, 5, (1, 1), (2, 2), (1, 1), 4, (10, 10, 14), (True, False, False), False],
- [3, 3, (2, 1), (1, 1), (1, 1), 7, (20, 20, 20), (False, False, False), False],
- [3, 3, (1, 1), (2, 2), (1, 1), 16, (10, 10, 14), (False, True, True), False],
- # Depth-wise convolution
- [3, 3, (1, 1), (1, 1), (1, 1), 20, (20, 20, 20), (False, False, True), True],
- [5, 5, (2, 2), (1, 1), (1, 1), 20, (20, 20, 20), (False, True, False), True],
- [3, 3, (2, 2), (2, 2), (1, 1), 14, (10, 10, 14), (True, False, False), True],
- [5, 5, (0, 0), (1, 1), (1, 1), 20, (20, 20, 20), (False, False, False), True],
- [3, 3, (1, 1), (2, 2), (1, 1), 14, (10, 10, 14), (False, True, True), True],
- ],
-)
-def test_conv2d(trial):
- Device.load("test_config.json")
-
- if skip_runtime_test():
- return
-
- device = Device()
- np.random.seed(0)
-
- dtype = "float32"
-
- (
- kernel_h,
- kernel_w,
- pad,
- stride,
- dilation,
- out_channels,
- shape,
- composite,
- is_depthwise,
- ) = trial
- shape = (1, *shape)
- if is_depthwise:
- groups = shape[3]
- else:
- groups = 1
- outputs = []
- inputs = {
- "a": tvm.nd.array(np.random.uniform(-128, 127, shape).astype(dtype)),
- }
-
- func, params = _get_model(
- shape,
- kernel_h,
- kernel_w,
- pad,
- stride,
- dilation,
- groups,
- dtype,
- out_channels,
- iter(inputs),
- has_pad=composite[0],
- has_bias=composite[1],
- has_activation=composite[2],
- )
- # Generate results for ACL conv2d and TVM native conv2d for comparison
- for acl in [False, True]:
- outputs.append(build_and_run(func, inputs, 1, params, device, enable_acl=acl)[0])
-
- config = {
- "shape": shape,
- "groups": groups,
- "kernel size": (kernel_h, kernel_w),
- "padding": pad,
- "stride": stride,
- "dilation": dilation,
- "out channels": out_channels,
- "composite operators (pad, bias, activation)": composite,
- }
- verify(outputs, atol=0.002, rtol=0.01, config=config)
-
-
-@pytest.mark.parametrize(
- "trial",
- [
- # Normal convolution
- [2, 2, (1, 1), (1, 1), (1, 1), 4, (10, 10, 14), (False, False, False), False],
- [2, 1, (2, 2), (1, 1), (1, 1), 7, (12, 15, 16), (False, False, True), False],
- [3, 3, (2, 1), (1, 1), (1, 1), 4, (10, 10, 14), (False, True, False), False],
- [3, 3, (1, 1), (1, 1), (1, 1), 16, (12, 15, 16), (False, False, False), False],
- [5, 5, (1, 1), (2, 2), (1, 1), 4, (10, 10, 14), (True, False, False), False],
- [1, 3, (1, 1), (1, 1), (1, 1), 7, (20, 20, 20), (False, False, True), False],
- [2, 2, (2, 2), (1, 1), (1, 1), 4, (20, 20, 20), (False, True, False), False],
- [5, 5, (1, 1), (2, 2), (1, 1), 4, (10, 10, 14), (True, False, False), False],
- [3, 3, (2, 1), (1, 1), (1, 1), 7, (20, 20, 20), (False, False, False), False],
- [3, 3, (1, 1), (2, 2), (1, 1), 16, (10, 10, 14), (False, True, True), False],
- # Depth-wise convolution
- [3, 3, (1, 1), (1, 1), (1, 1), 20, (20, 20, 20), (False, False, True), True],
- [5, 5, (2, 2), (1, 1), (1, 1), 20, (20, 20, 20), (False, True, False), True],
- [3, 3, (2, 2), (2, 2), (1, 1), 14, (10, 10, 14), (True, False, False), True],
- [5, 5, (0, 0), (1, 1), (1, 1), 20, (20, 20, 20), (False, False, False), True],
- [3, 3, (1, 1), (2, 2), (1, 1), 14, (10, 10, 14), (False, True, True), True],
- ],
-)
-def test_codegen_conv2d(trial):
- if skip_codegen_test():
- return
-
- dtype = "float32"
-
- (
- kernel_h,
- kernel_w,
- pad,
- stride,
- dilation,
- out_channels,
- shape,
- composite,
- is_depthwise,
- ) = trial
- shape = (1, *shape)
- if is_depthwise:
- groups = shape[3]
- else:
- groups = 1
- inputs = {"a"}
-
- args = (shape, kernel_h, kernel_w, pad, stride, dilation, groups, dtype, out_channels)
-
- func, params = _get_model(
- *args,
- var_names=iter(inputs),
- has_pad=composite[0],
- has_bias=composite[1],
- has_activation=composite[2],
- )
- exp_codegen = _get_expected_codegen(*args, has_bias=composite[1], has_activation=composite[2])
- verify_codegen(func, exp_codegen, 1)
-
-
-@pytest.mark.parametrize(
- "trial",
- [
- # Normal convolution
- [2, 2, (1, 1), (1, 1), (1, 1), 4, (10, 10, 14), (False, False, False), False],
- [2, 1, (2, 2), (1, 1), (1, 1), 7, (12, 15, 16), (False, False, True), False],
- [3, 3, (2, 1), (1, 1), (1, 1), 4, (10, 10, 14), (False, True, False), False],
- [3, 3, (1, 1), (1, 1), (1, 1), 16, (12, 15, 16), (False, False, False), False],
- [5, 5, (1, 1), (2, 2), (1, 1), 4, (10, 10, 14), (True, False, False), False],
- [1, 3, (1, 1), (1, 1), (1, 1), 7, (20, 20, 20), (False, False, True), False],
- [2, 2, (2, 2), (1, 1), (1, 1), 4, (20, 20, 20), (False, True, False), False],
- [5, 5, (1, 1), (2, 2), (1, 1), 4, (10, 10, 14), (True, False, False), False],
- [3, 3, (2, 1), (1, 1), (1, 1), 7, (20, 20, 20), (False, False, False), False],
- [3, 3, (1, 1), (2, 2), (1, 1), 16, (10, 10, 14), (False, True, True), False],
- # Depth-wise convolution
- [3, 3, (1, 1), (1, 1), (1, 1), 20, (20, 20, 20), (False, False, True), True],
- [5, 5, (2, 2), (1, 1), (1, 1), 20, (20, 20, 20), (False, True, False), True],
- [3, 3, (2, 2), (2, 2), (1, 1), 14, (10, 10, 14), (True, False, False), True],
- [5, 5, (0, 0), (1, 1), (1, 1), 20, (20, 20, 20), (False, False, False), True],
- [3, 3, (1, 1), (2, 2), (1, 1), 14, (10, 10, 14), (False, True, True), True],
- ],
-)
-@pytest.mark.parametrize("dtype", QNN_DTYPES)
-def test_qnn_conv2d(trial, dtype):
- Device.load("test_config.json")
-
- if skip_runtime_test():
- return
-
- device = Device()
- np.random.seed(0)
-
- (
- kernel_h,
- kernel_w,
- pad,
- stride,
- dilation,
- out_channels,
- shape,
- composite,
- is_depthwise,
- ) = trial
- shape = (1, *shape)
- if is_depthwise:
- groups = shape[3]
- else:
- groups = 1
- outputs = []
- inputs = {"a": tvm.nd.array(np.random.uniform(0, 255, shape).astype(dtype))}
-
- input_zp = 100
- input_sc = 0.5
- kernel_zp = 25
- kernel_sc = 0.03
- output_zp, output_sc = _get_qnn_params(
- input_zp, input_sc, kernel_zp, kernel_sc, kernel_h, kernel_w, shape[3]
- )
-
- func, params = _get_qnn_model(
- shape,
- kernel_h,
- kernel_w,
- pad,
- stride,
- dilation,
- groups,
- dtype,
- out_channels,
- input_zp,
- input_sc,
- kernel_zp,
- kernel_sc,
- output_zp,
- output_sc,
- iter(inputs),
- has_pad=composite[0],
- has_bias=composite[1],
- has_activation=composite[2],
- )
- for acl in [False, True]:
- outputs.append(build_and_run(func, inputs, 1, params, device, enable_acl=acl)[0])
-
- config = {
- "shape": shape,
- "groups": groups,
- "kernel size": (kernel_h, kernel_w),
- "padding": pad,
- "stride": stride,
- "dilation": dilation,
- "out channels": out_channels,
- "composite operators (pad, bias, activation)": composite,
- "input scale": input_sc,
- "input zero point": input_zp,
- "kernel scale": kernel_sc,
- "kernel zero point": kernel_zp,
- "output scale": output_sc,
- "output zero point": output_zp,
- }
-
- atol = 2 if is_depthwise else 1
- verify(outputs, atol=atol, rtol=0, config=config, verify_saturation=True)
-
-
-@pytest.mark.parametrize("dtype", QNN_DTYPES)
-@pytest.mark.parametrize(
- "trial",
- [
- # Normal convolution
- [2, 2, (1, 1), (1, 1), (1, 1), 4, (10, 10, 14), (False, False, False), False],
- [2, 1, (2, 2), (1, 1), (1, 1), 7, (12, 15, 16), (False, False, True), False],
- [3, 3, (2, 1), (1, 1), (1, 1), 4, (10, 10, 14), (False, True, False), False],
- [3, 3, (1, 1), (1, 1), (1, 1), 16, (12, 15, 16), (False, False, False), False],
- [5, 5, (1, 1), (2, 2), (1, 1), 4, (10, 10, 14), (True, False, False), False],
- [1, 3, (1, 1), (1, 1), (1, 1), 7, (20, 20, 20), (False, False, True), False],
- [2, 2, (2, 2), (1, 1), (1, 1), 4, (20, 20, 20), (False, True, False), False],
- [5, 5, (1, 1), (2, 2), (1, 1), 4, (10, 10, 14), (True, False, False), False],
- [3, 3, (2, 1), (1, 1), (1, 1), 7, (20, 20, 20), (False, False, False), False],
- [3, 3, (1, 1), (2, 2), (1, 1), 16, (10, 10, 14), (False, True, True), False],
- # Depth-wise convolution
- [3, 3, (1, 1), (1, 1), (1, 1), 20, (20, 20, 20), (False, False, True), True],
- [5, 5, (2, 2), (1, 1), (1, 1), 20, (20, 20, 20), (False, True, False), True],
- [3, 3, (2, 2), (2, 2), (1, 1), 14, (10, 10, 14), (True, False, False), True],
- [5, 5, (0, 0), (1, 1), (1, 1), 20, (20, 20, 20), (False, False, False), True],
- [3, 3, (1, 1), (2, 2), (1, 1), 14, (10, 10, 14), (False, True, True), True],
- ],
-)
-def test_codegen_qnn_conv2d(trial, dtype):
- if skip_codegen_test():
- return
-
- (
- kernel_h,
- kernel_w,
- pad,
- stride,
- dilation,
- out_channels,
- shape,
- composite,
- is_depthwise,
- ) = trial
- shape = (1, *shape)
- if is_depthwise:
- groups = shape[3]
- else:
- groups = 1
- inputs = {"a"}
-
- input_zp = 100
- input_sc = 0.5
- kernel_zp = 25
- kernel_sc = 0.03
- output_zp, output_sc = _get_qnn_params(
- input_zp, input_sc, kernel_zp, kernel_sc, kernel_h, kernel_w, shape[3]
- )
-
- args = (shape, kernel_h, kernel_w, pad, stride, dilation, groups, dtype, out_channels)
-
- func, params = _get_qnn_model(
- *args,
- input_zp=input_zp,
- input_sc=input_sc,
- kernel_zp=kernel_zp,
- kernel_sc=kernel_sc,
- output_zp=output_zp,
- output_sc=output_sc,
- var_names=iter(inputs),
- has_pad=composite[0],
- has_bias=composite[1],
- has_activation=composite[2],
- )
- exp_codegen = _get_expected_codegen(*args, has_bias=composite[1], has_activation=composite[2])
- verify_codegen(func, exp_codegen, 1)
-
-
-@pytest.mark.parametrize(
- "param",
- ["kernel_sc", "kernel_zp"],
-)
-def test_codegen_qnn_conv2d_per_channel_quantization(param):
- if skip_codegen_test():
- return
-
- dtype = "int8"
- kernel_h = 2
- kernel_w = 2
- pad = (1, 1)
- stride = (1, 1)
- dilation = (1, 1)
- out_channels = 4
- shape = (1, 10, 10, 14)
- composite = (False, False, False)
- groups = 1
- inputs = {"a"}
-
- qnn_params = {
- "input_zp": 1,
- "input_sc": 1,
- "kernel_zp": 1,
- "kernel_sc": 1,
- "output_zp": 1,
- "output_sc": 1,
- }
- qnn_params[param] = [1, 1, 1, 1]
-
- args = (shape, kernel_h, kernel_w, pad, stride, dilation, groups, dtype, out_channels)
-
- func, params = _get_qnn_model(
- *args,
- input_zp=qnn_params["input_zp"],
- input_sc=qnn_params["input_sc"],
- kernel_zp=qnn_params["kernel_zp"],
- kernel_sc=qnn_params["kernel_sc"],
- output_zp=qnn_params["output_zp"],
- output_sc=qnn_params["output_sc"],
- var_names=iter(inputs),
- has_pad=composite[0],
- has_bias=composite[1],
- has_activation=composite[2],
- )
- exp_codegen = _get_expected_codegen(*args, has_bias=composite[1], has_activation=composite[2])
- verify_codegen(func, exp_codegen, num_acl_modules=0, tvm_ops=2)
-
-
-if __name__ == "__main__":
- test_conv2d()
- test_qnn_conv2d()
- test_codegen_conv2d()
- test_codegen_qnn_conv2d()
- test_codegen_qnn_conv2d_per_channel_quantization()
diff --git a/tests/python/contrib/test_arm_compute_lib/test_dense.py b/tests/python/contrib/test_arm_compute_lib/test_dense.py
deleted file mode 100644
index 411f790f34..0000000000
--- a/tests/python/contrib/test_arm_compute_lib/test_dense.py
+++ /dev/null
@@ -1,430 +0,0 @@
-# Licensed to the Apache Software Foundation (ASF) under one
-# or more contributor license agreements. See the NOTICE file
-# distributed with this work for additional information
-# regarding copyright ownership. The ASF licenses this file
-# to you under the Apache License, Version 2.0 (the
-# "License"); you may not use this file except in compliance
-# with the License. You may obtain a copy of the License at
-#
-# http://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing,
-# software distributed under the License is distributed on an
-# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
-# KIND, either express or implied. See the License for the
-# specific language governing permissions and limitations
-# under the License.
-"""Arm Compute Library integration dense tests."""
-import numpy as np
-import pytest
-
-import tvm
-from tvm import relay
-from tvm import testing
-from test_arm_compute_lib.infrastructure import (
- Device,
- skip_runtime_test,
- skip_codegen_test,
- build_and_run,
- verify,
- verify_codegen,
-)
-
-
-def _get_model(shape, weight_shape, units, dtype, var_names, has_bias=False):
- """Return a model and any parameters it may have"""
- a = relay.var(next(var_names), shape=shape, dtype=dtype)
- w = tvm.nd.array(np.random.uniform(-128, 127, weight_shape).astype(dtype))
- weights = relay.const(w, dtype)
- out = relay.nn.dense(a, weights, units=units, out_dtype=dtype)
- params = {"w": w}
- if has_bias:
- b = tvm.nd.array(np.random.randint(-128, 127, weight_shape[0]).astype(dtype))
- biasc = relay.const(b, dtype)
- out = relay.nn.bias_add(out, biasc)
- params["b"] = b
- return out, params
-
-
-def _get_qnn_params(input_zp, input_sc, kernel_zp, kernel_sc, kernel_h, kernel_w):
- """Get output qnn parameters given input and kernel parameters."""
- input_max = input_sc * (255 - input_zp)
- input_min = -input_sc * input_zp
- kernel_max = kernel_sc * (255 - kernel_zp)
- kernel_min = -kernel_sc * kernel_zp
- output_limits = [
- kernel_max * kernel_h * kernel_w * input_max,
- kernel_min * kernel_h * kernel_w * input_max,
- kernel_min * kernel_h * kernel_w * input_min,
- kernel_max * kernel_h * kernel_w * input_min,
- ]
- output_max = max(output_limits)
- output_min = min(output_limits)
- output_sc = (output_max - output_min) / 255
- output_zp = -int(output_min / output_sc)
- return output_zp, output_sc
-
-
-def _get_qnn_model(
- shape,
- weight_shape,
- units,
- dtype,
- input_zp,
- input_sc,
- kernel_zp,
- kernel_sc,
- output_zp,
- output_sc,
- var_names,
- has_bias=False,
-):
- a = relay.var(next(var_names), shape=shape, dtype=dtype)
- w = tvm.nd.array(np.random.uniform(-128, 127, weight_shape).astype(dtype))
- weights = relay.const(w, dtype)
- out = relay.qnn.op.dense(
- a,
- weights,
- units=units,
- input_zero_point=relay.const(input_zp, "int32"),
- kernel_zero_point=relay.const(kernel_zp, "int32"),
- input_scale=relay.const(input_sc, "float32"),
- kernel_scale=relay.const(kernel_sc, "float32"),
- out_dtype="int32",
- )
- params = {"w": w}
- if has_bias:
- b = tvm.nd.array(np.random.randint(0, 255, weight_shape[0]).astype("int32"))
- biasc = relay.const(b, "int32")
- out = relay.nn.bias_add(out, biasc)
- params["b"] = b
- out = relay.qnn.op.requantize(
- out,
- relay.const(input_sc * kernel_sc, "float32"), # input scale
- relay.const(0, "int32"), # input zero point
- relay.const(output_sc, "float32"), # output scale
- relay.const(output_zp, "int32"), # output zero point
- out_dtype=dtype,
- )
- return out, params
-
-
-def _get_expected_codegen(shape, weight_shape, units, dtype, has_bias=False):
- output_shape = (shape[0], units)
- qnn_dtypes = ("uint8", "int8")
- out_dtype = "int32" if dtype in qnn_dtypes else "float32"
-
- node = {
- "op": "kernel",
- "name": "nn.dense",
- "inputs": [],
- "attrs": {
- "num_outputs": "1",
- "out_dtype": [[out_dtype]],
- "shape": [[list(output_shape)]],
- "dtype": [[dtype]],
- "units": [[str(units)]],
- },
- }
-
- inputs = [
- {"op": "input", "name": "", "attrs": {"shape": [[list(shape)]], "dtype": [[str(dtype)]]}},
- {
- "op": "const",
- "name": "",
- "attrs": {"shape": [[list(weight_shape)]], "dtype": [[str(dtype)]]},
- },
- ]
-
- # qnn.dense params, input and kernel
- if dtype in qnn_dtypes:
- node["name"] = "qnn.dense"
- for param_dtype in ["int32", "float32"]:
- for _ in range(2):
- inputs.append(
- {
- "op": "const",
- "name": "",
- "attrs": {"shape": [[[]]], "dtype": [[param_dtype]]},
- }
- )
-
- if has_bias:
- bias_dtype = "int32" if dtype in qnn_dtypes else "float32"
- bias_shape = [1, weight_shape[0]] if weight_shape[0] != 1 else [weight_shape[0]]
- inputs.append(
- {
- "op": "const",
- "name": "",
- "attrs": {"shape": [[bias_shape]], "dtype": [[bias_dtype]]},
- }
- )
-
- # qnn.dense params, output
- if dtype in qnn_dtypes:
- for param_dtype in ["float32", "int32"]:
- inputs.append(
- {"op": "const", "name": "", "attrs": {"shape": [[[]]], "dtype": [[param_dtype]]}}
- )
-
- input_idx = 0
- for _ in range(len(inputs)):
- node["inputs"].append([input_idx, 0, 0])
- input_idx += 1
- node["attrs"]["num_inputs"] = str(len(inputs))
- inputs.append(node)
- return inputs
-
-
-def test_dense():
- Device.load("test_config.json")
-
- if skip_runtime_test():
- return
-
- device = Device()
- np.random.seed(0)
- dtype = "float32"
- trials = [
- [(1, 128), (16, 128), 16, True],
- [(1, 128), (16, 128), 16, False],
- [(32, 32), (32, 32), 32, True],
- [(32, 32), (32, 32), 32, False],
- [(1, 64), (1, 64), 1, True],
- [(1, 64), (1, 64), 1, False],
- [(11, 2), (2, 2), 2, True],
- [(11, 2), (2, 2), 2, False],
- ]
- for shape, weight_shape, units, composite in trials:
- outputs = []
- inputs = {"a": tvm.nd.array(np.random.uniform(-128, 127, shape).astype(dtype))}
- func, params = _get_model(
- shape, weight_shape, units, dtype, var_names=iter(inputs), has_bias=composite
- )
- for acl in [False, True]:
- outputs.append(
- build_and_run(
- func,
- inputs,
- 1,
- params,
- device,
- enable_acl=acl,
- )[0]
- )
- config = {
- "shape": shape,
- "weight_shape": weight_shape,
- "units": units,
- "dtype": dtype,
- "composite operators (bias)": composite,
- }
- verify(outputs, atol=0.001, rtol=0.01, config=config)
-
-
-def test_codegen_dense():
- if skip_codegen_test():
- return
-
- np.random.seed(0)
- dtype = "float32"
- trials = [
- [(1, 128), (16, 128), 16, True],
- [(1, 128), (16, 128), 16, False],
- [(32, 32), (32, 32), 32, True],
- [(32, 32), (32, 32), 32, False],
- [(1, 64), (1, 64), 1, True],
- [(1, 64), (1, 64), 1, False],
- [(11, 2), (2, 2), 2, True],
- [(11, 2), (2, 2), 2, False],
- ]
- for shape, weight_shape, units, composite in trials:
- inputs = {"a"}
-
- args = (shape, weight_shape, units, dtype)
-
- func, params = _get_model(*args, var_names=iter(inputs), has_bias=composite)
- exp_codegen = _get_expected_codegen(*args, has_bias=composite)
- verify_codegen(func, exp_codegen)
-
-
-@pytest.mark.parametrize(
- "dtype,min_range,max_range",
- [
- ("uint8", 0, 255),
- ("int8", -127, 128),
- ],
-)
-def test_qnn_dense(dtype, min_range, max_range):
- Device.load("test_config.json")
-
- if skip_runtime_test():
- return
-
- device = Device()
- np.random.seed(0)
-
- trials = [
- [(1, 2), (2, 2), 2, True],
- [(1, 2), (2, 2), 2, False],
- [(4, 4), (4, 4), 4, True],
- [(4, 4), (4, 4), 4, False],
- [(16, 16), (4, 16), 4, True],
- [(16, 16), (4, 16), 4, False],
- [(1, 128), (16, 128), 16, True],
- [(1, 128), (16, 128), 16, False],
- [(32, 32), (32, 32), 32, True],
- [(32, 32), (32, 32), 32, False],
- [(1, 64), (1, 64), 1, True],
- [(1, 64), (1, 64), 1, False],
- ]
- for shape, weight_shape, units, composite in trials:
- outputs = []
- inputs = {"a": tvm.nd.array(np.random.uniform(min_range, max_range, shape).astype(dtype))}
- input_zp = 100
- input_sc = 0.5
- kernel_zp = 50
- kernel_sc = 0.03
- output_zp, output_sc = _get_qnn_params(
- input_zp, input_sc, kernel_zp, kernel_sc, weight_shape[0], weight_shape[1]
- )
-
- func, params = _get_qnn_model(
- shape,
- weight_shape,
- units,
- dtype,
- input_zp,
- input_sc,
- kernel_zp,
- kernel_sc,
- output_zp,
- output_sc,
- var_names=iter(inputs),
- has_bias=composite,
- )
-
- for acl in [False, True]:
- outputs.append(
- build_and_run(
- func,
- inputs,
- 1,
- params,
- device,
- enable_acl=acl,
- )[0]
- )
-
- config = {
- "shape": shape,
- "weight_shape": weight_shape,
- "units": units,
- "dtype": dtype,
- "composite operators (bias)": composite,
- "input scale": input_sc,
- "input zero point": input_zp,
- "kernel scale": kernel_sc,
- "kernel zero point": kernel_zp,
- "output scale": output_sc,
- "output zero point": output_zp,
- }
- verify(outputs, atol=1, rtol=0, config=config, verify_saturation=True)
-
-
-@pytest.mark.parametrize("dtype", ["uint8", "int8"])
-def test_codegen_qnn_dense(dtype):
- if skip_codegen_test():
- return
-
- np.random.seed(0)
-
- trials = [
- [(1, 2), (2, 2), 2, True],
- [(1, 2), (2, 2), 2, False],
- [(4, 4), (4, 4), 4, True],
- [(4, 4), (4, 4), 4, False],
- [(16, 16), (4, 16), 4, True],
- [(16, 16), (4, 16), 4, False],
- [(1, 128), (16, 128), 16, True],
- [(1, 128), (16, 128), 16, False],
- [(32, 32), (32, 32), 32, True],
- [(32, 32), (32, 32), 32, False],
- [(1, 64), (1, 64), 1, True],
- [(1, 64), (1, 64), 1, False],
- ]
- for shape, weight_shape, units, composite in trials:
- inputs = {"a"}
- args = (shape, weight_shape, units, dtype)
-
- input_zp = 100
- input_sc = 0.5
- kernel_zp = 25
- kernel_sc = 0.03
- output_zp, output_sc = _get_qnn_params(
- input_zp, input_sc, kernel_zp, kernel_sc, weight_shape[0], weight_shape[1]
- )
-
- func, params = _get_qnn_model(
- *args,
- var_names=iter(inputs),
- input_zp=input_zp,
- input_sc=input_sc,
- kernel_zp=kernel_zp,
- kernel_sc=kernel_sc,
- output_zp=output_zp,
- output_sc=output_sc,
- has_bias=composite,
- )
- exp_codegen = _get_expected_codegen(*args, has_bias=composite)
- verify_codegen(func, exp_codegen)
-
-
-@pytest.mark.parametrize(
- "param",
- ["kernel_sc", "kernel_zp"],
-)
-def test_codegen_qnn_dense_per_channel_quantization(param):
- if skip_codegen_test():
- return
-
- np.random.seed(0)
- dtype = "int8"
- shape = (1, 2)
- weight_shape = (2, 2)
- units = 2
- composite = True
- inputs = {"a"}
- args = (shape, weight_shape, units, dtype)
-
- qnn_params = {
- "input_zp": 1,
- "input_sc": 1,
- "kernel_zp": 1,
- "kernel_sc": 1,
- "output_zp": 1,
- "output_sc": 1,
- }
- qnn_params[param] = [1, 1]
-
- func, _ = _get_qnn_model(
- *args,
- var_names=iter(inputs),
- input_zp=qnn_params["input_zp"],
- input_sc=qnn_params["input_sc"],
- kernel_zp=qnn_params["kernel_zp"],
- kernel_sc=qnn_params["kernel_sc"],
- output_zp=qnn_params["output_zp"],
- output_sc=qnn_params["output_sc"],
- has_bias=composite,
- )
- exp_codegen = _get_expected_codegen(*args, has_bias=composite)
- verify_codegen(func, exp_codegen, num_acl_modules=0, tvm_ops=3)
-
-
-if __name__ == "__main__":
- test_dense()
- test_qnn_dense()
- test_codegen_dense()
- test_codegen_qnn_dense()
- test_codegen_qnn_dense_per_channel_quantization()
diff --git a/tests/python/contrib/test_arm_compute_lib/test_maximum.py b/tests/python/contrib/test_arm_compute_lib/test_maximum.py
deleted file mode 100644
index 1942d1e213..0000000000
--- a/tests/python/contrib/test_arm_compute_lib/test_maximum.py
+++ /dev/null
@@ -1,105 +0,0 @@
-# Licensed to the Apache Software Foundation (ASF) under one
-# or more contributor license agreements. See the NOTICE file
-# distributed with this work for additional information
-# regarding copyright ownership. The ASF licenses this file
-# to you under the Apache License, Version 2.0 (the
-# "License"); you may not use this file except in compliance
-# with the License. You may obtain a copy of the License at
-#
-# http://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing,
-# software distributed under the License is distributed on an
-# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
-# KIND, either express or implied. See the License for the
-# specific language governing permissions and limitations
-# under the License.
-"""Arm Compute Library integration reshape tests."""
-
-import numpy as np
-
-import tvm
-from tvm import relay
-from tvm import testing
-
-from .infrastructure import (
- skip_runtime_test,
- skip_codegen_test,
- build_and_run,
- verify,
- verify_codegen,
-)
-from .infrastructure import Device
-
-
-def _get_model(input_shape, dtype, var_names):
- """Return a model and any parameters it may have."""
- a = relay.var(next(var_names), shape=input_shape, dtype=dtype)
- b = relay.var(next(var_names), shape=input_shape, dtype=dtype)
- max = relay.maximum(a, b)
- return max
-
-
-def _get_expected_codegen(shape, dtype):
- node = {
- "op": "kernel",
- "name": "maximum",
- "inputs": [[0, 0, 0], [1, 0, 0]],
- "attrs": {
- "num_inputs": "2",
- "num_outputs": "1",
- "shape": [[list(shape)]],
- "dtype": [[dtype]],
- },
- }
-
- inputs = [
- {"op": "input", "name": "", "attrs": {"shape": [[list(shape)]], "dtype": [[dtype]]}},
- {"op": "input", "name": "", "attrs": {"shape": [[list(shape)]], "dtype": [[dtype]]}},
- ]
- inputs.append(node)
- return inputs
-
-
-def test_maximum():
- Device.load("test_config.json")
-
- if skip_runtime_test():
- return
-
- device = Device()
- np.random.seed(0)
-
- for dtype, low, high, atol, rtol in [
- ("float32", -127, 128, 0.001, 0.001),
- ("float32", -1, 1, 0.001, 0.001),
- ]:
- inputs = {
- "a": tvm.nd.array(np.random.uniform(low, high, (100, 100)).astype(dtype)),
- "b": tvm.nd.array(np.random.uniform(low, high, (100, 100)).astype(dtype)),
- }
- outputs = []
- func = _get_model(inputs["a"].shape, dtype, iter(inputs))
-
- for acl in [False, True]:
- outputs.append(build_and_run(func, inputs, 1, None, device, enable_acl=acl)[0])
-
- verify(outputs, atol=1e-7, rtol=1e-7)
-
-
-def test_codegen_maximum():
- if skip_codegen_test():
- return
-
- shape = (100, 100)
- inputs = {"a", "b"}
- for dtype in ["float32"]:
- args = (shape, dtype)
- func = _get_model(*args, iter(inputs))
- exp_codegen = _get_expected_codegen(*args)
- verify_codegen(func, exp_codegen, 1)
-
-
-if __name__ == "__main__":
- test_maximum()
- test_codegen_maximum()
diff --git a/tests/python/contrib/test_arm_compute_lib/test_network.py b/tests/python/contrib/test_arm_compute_lib/test_network.py
deleted file mode 100644
index 8c6302abf8..0000000000
--- a/tests/python/contrib/test_arm_compute_lib/test_network.py
+++ /dev/null
@@ -1,231 +0,0 @@
-# Licensed to the Apache Software Foundation (ASF) under one
-# or more contributor license agreements. See the NOTICE file
-# distributed with this work for additional information
-# regarding copyright ownership. The ASF licenses this file
-# to you under the Apache License, Version 2.0 (the
-# "License"); you may not use this file except in compliance
-# with the License. You may obtain a copy of the License at
-#
-# http://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing,
-# software distributed under the License is distributed on an
-# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
-# KIND, either express or implied. See the License for the
-# specific language governing permissions and limitations
-# under the License.
-"""Arm Compute Library network tests."""
-
-from packaging.version import parse
-
-import numpy as np
-import pytest
-from tvm import relay
-
-from test_arm_compute_lib.infrastructure import Device, skip_runtime_test, build_and_run, verify
-
-
-def _build_and_run_network(mod, params, inputs, device, tvm_ops, acl_partitions, atol, rtol):
- """Helper function to build and run a network."""
- data = {}
- np.random.seed(0)
-
- for name, (shape, dtype) in inputs.items():
- if dtype == "uint8":
- low, high = 0, 255
- else:
- low, high = -127, 128
- data[name] = np.random.uniform(low, high, shape).astype(dtype)
-
- outputs = []
- for acl in [False, True]:
- outputs.append(
- build_and_run(
- mod,
- data,
- 1,
- params,
- device,
- enable_acl=acl,
- tvm_ops=tvm_ops,
- acl_partitions=acl_partitions,
- )[0]
- )
- verify(outputs, atol=atol, rtol=rtol, verify_saturation=False)
-
-
-def _get_tflite_model(tflite_model_path, inputs_dict):
- """Convert TFlite graph to relay."""
- try:
- import tflite.Model
- except ImportError:
- pytest.skip("Missing Tflite support")
-
- with open(tflite_model_path, "rb") as f:
- tflite_model_buffer = f.read()
-
- try:
- tflite_model = tflite.Model.Model.GetRootAsModel(tflite_model_buffer, 0)
- except AttributeError:
- tflite_model = tflite.Model.GetRootAsModel(tflite_model_buffer, 0)
- shape_dict = {}
- dtype_dict = {}
- for input in inputs_dict:
- input_shape, input_dtype = inputs_dict[input]
- shape_dict[input] = input_shape
- dtype_dict[input] = input_dtype
-
- return relay.frontend.from_tflite(tflite_model, shape_dict=shape_dict, dtype_dict=dtype_dict)
-
-
-def _get_keras_model(keras_model, inputs_dict):
- """Convert Keras graph to relay."""
- inputs = {}
- for name, (shape, _) in inputs_dict.items():
- inputs[keras_model.input_names[0]] = shape
- return relay.frontend.from_keras(keras_model, inputs, layout="NHWC")
-
-
-def test_vgg16():
- Device.load("test_config.json")
-
- if skip_runtime_test():
- return
-
- device = Device()
-
- def get_model():
- try:
- from keras.applications import VGG16
- except ImportError:
- pytest.skip("Missing Keras Package")
-
- vgg16 = VGG16(include_top=True, weights="imagenet", input_shape=(224, 224, 3), classes=1000)
- inputs = {vgg16.input_names[0]: ((1, 224, 224, 3), "float32")}
- mod, params = _get_keras_model(vgg16, inputs)
- return mod, params, inputs
-
- _build_and_run_network(
- *get_model(),
- device=device,
- tvm_ops=4,
- acl_partitions=21,
- atol=0.002,
- rtol=0.01,
- )
-
-
-def test_mobilenet():
- keras = pytest.importorskip("keras")
- Device.load("test_config.json")
-
- if skip_runtime_test():
- return
-
- device = Device()
-
- def get_model():
- try:
- from keras.applications import MobileNet
- except ImportError:
- pytest.skip("Missing keras module")
-
- mobilenet = MobileNet(
- include_top=True, weights="imagenet", input_shape=(224, 224, 3), classes=1000
- )
- inputs = {mobilenet.input_names[0]: ((1, 224, 224, 3), "float32")}
- mod, params = _get_keras_model(mobilenet, inputs)
- return mod, params, inputs
-
- if parse(keras.__version__) < parse("2.9"):
- # This can be removed after we migrate to TF/Keras >= 2.9
- expected_tvm_ops = 56
- expected_acl_partitions = 31
- else:
- # In Keras >= 2.7, one reshape operator was removed
- # from the MobileNet model, so it impacted this test
- # which now needs to be reduce in by 1
- # The change in Keras is `b6abfaed1326e3c`
- expected_tvm_ops = 55
- expected_acl_partitions = 30
-
- _build_and_run_network(
- *get_model(),
- device=device,
- tvm_ops=expected_tvm_ops,
- acl_partitions=expected_acl_partitions,
- atol=0.002,
- rtol=0.01,
- )
-
-
-def test_quantized_mobilenet():
- Device.load("test_config.json")
-
- if skip_runtime_test():
- return
-
- try:
- import tvm.relay.testing.tf as tf_testing
- except ImportError:
- pytest.skip("Missing Tflite support")
-
- device = Device()
-
- def get_model():
- model_path = tf_testing.get_workload_official(
- "https://storage.googleapis.com/download.tensorflow.org/"
- "models/mobilenet_v1_2018_08_02/mobilenet_v1_1.0_224_quant.tgz",
- "mobilenet_v1_1.0_224_quant.tflite",
- )
- inputs = {"input": ((1, 224, 224, 3), "uint8")}
- mod, params = _get_tflite_model(model_path, inputs_dict=inputs)
- return mod, params, inputs
-
- _build_and_run_network(
- *get_model(),
- device=device,
- tvm_ops=3,
- acl_partitions=30,
- atol=10,
- rtol=0,
- )
-
-
-def test_squeezenet():
- Device.load("test_config.json")
-
- if skip_runtime_test():
- return
-
- try:
- import tvm.relay.testing.tf as tf_testing
- except ImportError:
- pytest.skip("Missing TF Support")
-
- device = Device()
-
- def get_model():
- model_path = tf_testing.get_workload_official(
- "https://storage.googleapis.com/download.tensorflow.org/models/tflite/model_zoo/upload_20180427/squeezenet_2018_04_27.tgz",
- "squeezenet.tflite",
- )
- inputs = {"Placeholder": ((1, 224, 224, 3), "float32")}
- mod, params = _get_tflite_model(model_path, inputs_dict=inputs)
- return mod, params, inputs
-
- _build_and_run_network(
- *get_model(),
- device=device,
- tvm_ops=9,
- acl_partitions=31,
- atol=8,
- rtol=0,
- )
-
-
-if __name__ == "__main__":
- test_vgg16()
- test_mobilenet()
- test_quantized_mobilenet()
- test_squeezenet()
diff --git a/tests/python/contrib/test_arm_compute_lib/test_pooling.py b/tests/python/contrib/test_arm_compute_lib/test_pooling.py
deleted file mode 100644
index f08fa0059d..0000000000
--- a/tests/python/contrib/test_arm_compute_lib/test_pooling.py
+++ /dev/null
@@ -1,419 +0,0 @@
-# Licensed to the Apache Software Foundation (ASF) under one
-# or more contributor license agreements. See the NOTICE file
-# distributed with this work for additional information
-# regarding copyright ownership. The ASF licenses this file
-# to you under the Apache License, Version 2.0 (the
-# "License"); you may not use this file except in compliance
-# with the License. You may obtain a copy of the License at
-#
-# http://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing,
-# software distributed under the License is distributed on an
-# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
-# KIND, either express or implied. See the License for the
-# specific language governing permissions and limitations
-# under the License.
-"""Arm Compute Library integration pooling tests."""
-import numpy as np
-import pytest
-
-import tvm
-from tvm import relay, testing
-
-from test_arm_compute_lib.infrastructure import (
- Device,
- build_and_run,
- skip_codegen_test,
- skip_runtime_test,
- verify,
- verify_codegen,
-)
-
-
-def _calculate_output_shape(shape, sizes, padding, strides, dilation):
- """Calculate pooling output shape."""
- height_receptive_field = (sizes[0] - 1) * dilation[0] + 1
- width_receptive_field = (sizes[1] - 1) * dilation[1] + 1
- output_height = ((shape[1] - height_receptive_field + padding[0] + padding[2]) / strides[0]) + 1
- output_width = ((shape[2] - width_receptive_field + padding[1] + padding[3]) / strides[1]) + 1
- return 1, int(output_height), int(output_width), shape[3]
-
-
-def _get_pooling_model(
- shape, dtype, typef, sizes, strides, dilation, padding, ceil_mode, count_include_pad, var_names
-):
- """Return a model and any parameters it may have."""
- if len(padding) == 2:
- padding = (padding[0], padding[1], padding[0], padding[1])
- out = relay.var(next(var_names), shape=shape, dtype=dtype)
- qnn_dtypes = ("uint8", "int8")
-
- if typef == "nn.max_pool2d":
- out = relay.nn.max_pool2d(
- out,
- pool_size=sizes,
- strides=strides,
- dilation=dilation,
- padding=padding,
- ceil_mode=ceil_mode,
- layout="NHWC",
- )
- elif typef == "nn.avg_pool2d":
- if dtype in qnn_dtypes:
- out = relay.cast(out, "int32")
- out = relay.nn.avg_pool2d(
- out,
- pool_size=sizes,
- strides=strides,
- dilation=dilation,
- padding=padding,
- ceil_mode=ceil_mode,
- count_include_pad=count_include_pad,
- layout="NHWC",
- )
- if dtype in qnn_dtypes:
- out = relay.cast(out, dtype)
- elif typef == "nn.l2_pool2d":
- out = relay.power(out, relay.const(2.0))
- out = relay.nn.avg_pool2d(
- out,
- pool_size=sizes,
- strides=strides,
- padding=padding,
- ceil_mode=ceil_mode,
- count_include_pad=count_include_pad,
- layout="NHWC",
- )
- out = relay.sqrt(out)
- else:
- raise ValueError("Function not supported")
-
- return out
-
-
-def _get_global_pooling_model(shape, dtype, typef, var_names):
- """Return a model and any parameters it may have."""
- out = relay.var(next(var_names), shape=shape, dtype=dtype)
- qnn_dtypes = ("uint8", "int8")
-
- if typef == "nn.global_max_pool2d":
- out = relay.nn.global_max_pool2d(out, layout="NHWC")
- elif typef == "nn.global_avg_pool2d":
- if dtype in qnn_dtypes:
- out = relay.cast(out, "int32")
- out = relay.nn.global_avg_pool2d(out, layout="NHWC")
- if dtype in qnn_dtypes:
- out = relay.cast(out, dtype)
- else:
- raise ValueError("Function not supported")
-
- return out
-
-
-def _get_expected_pooling_codegen(
- shape, dtype, typef, sizes, strides, dilation, padding, ceil_mode, count_include_pad
-):
- if len(padding) == 2:
- padding = (padding[0], padding[1], padding[0], padding[1])
- output_shape = _calculate_output_shape(shape, sizes, padding, strides, dilation)
-
- node = {
- "op": "kernel",
- "name": typef,
- "inputs": [[0, 0, 0]],
- "attrs": {
- "num_inputs": "1",
- "num_outputs": "1",
- "layout": [["NHWC"]],
- "out_layout": [[""]],
- "shape": [[list(output_shape)]],
- "dtype": [[dtype]],
- "padding": [[str(p) for p in padding]],
- "strides": [[str(s) for s in strides]],
- "dilation": [[str(d) for d in dilation]],
- "pool_size": [[str(s) for s in sizes]],
- "ceil_mode": [[str(1 if ceil_mode else 0)]],
- },
- }
-
- if typef == "nn.avg_pool2d" or typef == "nn.l2_pool2d":
- node["attrs"]["count_include_pad"] = [["1" if count_include_pad else "0"]]
-
- input = {"op": "input", "name": "", "attrs": {"shape": [[list(shape)]], "dtype": [[dtype]]}}
- return [input, node]
-
-
-def _get_expected_global_pooling_codegen(shape, dtype, typef):
- node = {
- "op": "kernel",
- "name": typef,
- "inputs": [[0, 0, 0]],
- "attrs": {
- "num_inputs": "1",
- "num_outputs": "1",
- "layout": [["NHWC"]],
- "out_layout": [[""]],
- "shape": [[[1, 1, 1, shape[3]]]],
- "dtype": [[dtype]],
- },
- }
-
- input = {"op": "input", "name": "", "attrs": {"shape": [[list(shape)]], "dtype": [[dtype]]}}
- return [input, node]
-
-
-def _get_low_high_atol_rtol(dtype):
- if dtype == "float32":
- low, high, atol, rtol = (-127, 128, 0.001, 0.001)
- elif dtype == "uint8":
- low, high, atol, rtol = (0, 255, 1, 0)
- elif dtype == "int8":
- low, high, atol, rtol = (-127, 128, 1, 0)
- else:
- pytest.fail(f"dtype not expected: {dtype}")
-
- return low, high, atol, rtol
-
-
-# fmt: off
-@pytest.mark.parametrize(
- "typef,dtype,size,stride,dilation,pad,ceil_mode,count_include_pad,input_shape,expected_ops",
- [
- ("nn.max_pool2d", "float32", (3, 3), (2, 2), (1, 1), (0, 0), False, False, (27, 27, 512), (0, 1),),
- ("nn.max_pool2d", "float32", (2, 2), (2, 2), (1, 1), (0, 0), False, True, (16, 16, 16), (0, 1),),
- ("nn.max_pool2d", "float32", (3, 3), (2, 2), (1, 1), (1, 1), True, True, (15, 15, 16), (0, 1),),
- ("nn.max_pool2d", "float32", (2, 2), (2, 2), (1, 1), (0, 1), False, False, (16, 16, 16), (0, 1),),
- ("nn.max_pool2d", "uint8", (3, 3), (2, 2), (1, 1), (0, 1), False, False, (16, 16, 16), (0, 1),),
- ("nn.max_pool2d", "uint8", (2, 2), (2, 2), (1, 1), (1, 1), True, True, (15, 15, 16), (0, 1),),
- ("nn.max_pool2d", "uint8", (2, 2), (2, 2), (3, 2), (1, 1), True, True, (15, 15, 16), (1, 0),),
- ("nn.max_pool2d", "int8", (3, 3), (2, 2), (1, 1), (0, 1), False, False, (16, 16, 16), (0, 1),),
- ("nn.max_pool2d", "int8", (2, 2), (2, 2), (1, 1), (1, 1), True, True, (15, 15, 16), (0, 1),),
- ("nn.max_pool2d", "int8", (2, 2), (2, 2), (3, 2), (1, 1), True, True, (15, 15, 16), (1, 0),),
- ("nn.avg_pool2d", "float32", (2, 2), (2, 2), (1, 1), (1, 1), False, False, (16, 16, 16), (0, 1),),
- ("nn.avg_pool2d", "float32", (2, 2), (2, 2), (1, 1), (0, 0), False, True, (16, 16, 16), (0, 1),),
- ("nn.avg_pool2d", "float32", (3, 3), (2, 2), (3, 2), (0, 1), True, False, (15, 15, 16), (1, 0),),
- # 20.05: "exclude_padding equal false is not supported for AVG Pooling with padding on quantized types"
- # ["nn.avg_pool2d", uint8_dtype, (2, 2), (2, 2), (1, 1), False, True, (16, 16, 16)],
- ("nn.avg_pool2d", "uint8", (3, 3), (2, 2), (1, 1), (0, 1), False, False, (16, 16, 16), (0, 1),),
- ("nn.avg_pool2d", "int8", (3, 3), (2, 2), (1, 1), (0, 1), False, False, (16, 16, 16), (0, 1),),
- ("nn.l2_pool2d", "float32", (2, 2), (2, 2), (1, 1), (0, 1), True, False, (16, 16, 16), (0, 1),),
- ("nn.l2_pool2d", "float32", (3, 3), (2, 2), (1, 1), (0, 0), False, False, (16, 16, 16), (0, 1),),
- ("nn.l2_pool2d", "float32", (2, 2), (2, 2), (1, 1), (1, 1), False, True, (15, 15, 16), (0, 1),),
-
- ],
-)
-# fmt: on
-def test_pooling(
- typef,
- dtype,
- size,
- stride,
- dilation,
- pad,
- ceil_mode,
- count_include_pad,
- input_shape,
- expected_ops,
-):
- Device.load("test_config.json")
-
- if skip_runtime_test():
- return
-
- device = Device()
- np.random.seed(0)
-
- low, high, atol, rtol = _get_low_high_atol_rtol(dtype)
- tvm_ops, acl_partitions = expected_ops
-
- shape = (1, *input_shape)
- outputs = []
- inputs = {
- "a": tvm.nd.array(np.random.uniform(low, high, shape).astype(dtype)),
- }
-
- func = _get_pooling_model(
- shape,
- dtype,
- typef,
- size,
- stride,
- dilation,
- pad,
- ceil_mode,
- count_include_pad,
- iter(inputs),
- )
-
- config = {
- "size": size,
- "stride": stride,
- "shape": shape,
- "pooling type": typef,
- "dtype": dtype,
- "padding": pad,
- "dilation": dilation,
- "ceil_mode": ceil_mode,
- "count_include_pad": count_include_pad,
- "inputs": inputs,
- }
- verify_saturation = True if dtype == "uint8" else False
- for acl in [False, True]:
- outputs.append(
- build_and_run(
- func,
- inputs,
- 1,
- None,
- device,
- enable_acl=acl,
- tvm_ops=tvm_ops,
- acl_partitions=acl_partitions,
- config=config,
- )[0]
- )
-
- verify(outputs, atol=atol, rtol=rtol, config=config, verify_saturation=verify_saturation)
-
-
-@pytest.mark.parametrize(
- "typef,dtype,input_shape",
- [
- ["nn.global_max_pool2d", "float32", (8, 8, 16)],
- ["nn.global_max_pool2d", "float32", (9, 9, 16)],
- ["nn.global_max_pool2d", "uint8", (8, 8, 16)],
- ["nn.global_max_pool2d", "uint8", (9, 9, 16)],
- ["nn.global_max_pool2d", "int8", (8, 8, 16)],
- ["nn.global_max_pool2d", "int8", (9, 9, 16)],
- ["nn.global_avg_pool2d", "float32", (8, 8, 16)],
- ["nn.global_avg_pool2d", "float32", (9, 9, 16)],
- ["nn.global_avg_pool2d", "uint8", (8, 8, 16)],
- ["nn.global_avg_pool2d", "uint8", (9, 9, 16)],
- ["nn.global_avg_pool2d", "int8", (8, 8, 16)],
- ["nn.global_avg_pool2d", "int8", (9, 9, 16)],
- ],
-)
-def test_global_pooling(typef, dtype, input_shape):
- Device.load("test_config.json")
-
- if skip_runtime_test():
- return
-
- device = Device()
- np.random.seed(0)
-
- low, high, rtol, atol = _get_low_high_atol_rtol(dtype)
-
- shape = (1, *input_shape)
- outputs = []
- inputs = {
- "a": tvm.nd.array(np.random.uniform(low, high, shape).astype(dtype)),
- }
-
- func = _get_global_pooling_model(shape, dtype, typef, iter(inputs))
- config = {
- "shape": shape,
- "pooling type": typef,
- "dtype": dtype,
- }
- verify_saturation = True if dtype in ("uint8", "int8") else False
-
- for acl in [False, True]:
- outputs.append(
- build_and_run(func, inputs, 1, None, device, enable_acl=acl, config=config)[0]
- )
-
- verify(outputs, atol=atol, rtol=rtol, config=config, verify_saturation=verify_saturation)
-
-
-# fmt: off
-@pytest.mark.parametrize(
- "typef,dtype,size,stride,dilation,pad,ceil_mode,count_include_pad,input_shape,expected_ops",
- [
- ("nn.max_pool2d", "float32", (2, 2), (2, 2), (1, 1), (0, 0), False, True, (16, 16, 16), (0, 1),),
- ("nn.max_pool2d", "float32", (3, 3), (2, 2), (1, 1), (1, 1), True, True, (15, 15, 16), (0, 1),),
- ("nn.max_pool2d", "float32", (2, 2), (2, 2), (1, 1), (0, 1), False, False, (16, 16, 16), (0, 1),),
- ("nn.max_pool2d", "uint8", (3, 3), (2, 2), (1, 1), (0, 1), False, False, (16, 16, 16), (0, 1),),
- ("nn.max_pool2d", "uint8", (2, 2), (2, 2), (1, 1), (1, 1), True, True, (15, 15, 16), (0, 1),),
- ("nn.max_pool2d", "uint8", (2, 2), (2, 2), (3, 2), (1, 1), True, True, (15, 15, 16), (1, 0),),
- ("nn.max_pool2d", "int8", (3, 3), (2, 2), (1, 1), (0, 1), False, False, (16, 16, 16), (0, 1),),
- ("nn.max_pool2d", "int8", (2, 2), (2, 2), (1, 1), (1, 1), True, True, (15, 15, 16), (0, 1),),
- ("nn.max_pool2d", "int8", (2, 2), (2, 2), (3, 2), (1, 1), True, True, (15, 15, 16), (1, 0),),
- ("nn.avg_pool2d", "float32", (2, 2), (2, 2), (1, 1), (1, 1), False, False, (16, 16, 16), (0, 1),),
- ("nn.avg_pool2d", "float32", (2, 2), (2, 2), (1, 1), (1, 1), False, False, (16, 16, 16), (0, 1),),
- ("nn.avg_pool2d", "float32", (2, 2), (2, 2), (1, 1), (0, 0), False, True, (16, 16, 16), (0, 1),),
- ("nn.avg_pool2d", "float32", (3, 3), (2, 2), (3, 2), (0, 1), True, False, (15, 15, 16), (1, 0),),
- ("nn.avg_pool2d", "uint8", (2, 2), (2, 2), (1, 1), (1, 1), False, True, (16, 16, 16), (0, 1),),
- ("nn.avg_pool2d", "uint8", (3, 3), (2, 2), (1, 1), (0, 1), False, False, (16, 16, 16), (0, 1),),
- ("nn.avg_pool2d", "int8", (2, 2), (2, 2), (1, 1), (1, 1), False, True, (16, 16, 16), (0, 1),),
- ("nn.avg_pool2d", "int8", (3, 3), (2, 2), (1, 1), (0, 1), False, False, (16, 16, 16), (0, 1),),
- ("nn.l2_pool2d", "float32", (2, 2), (2, 2), (1, 1), (0, 1), True, False, (15, 15, 16), (0, 1),),
- ("nn.l2_pool2d", "float32", (3, 3), (2, 2), (1, 1), (0, 0), False, False, (16, 16, 16), (0, 1),),
- ("nn.l2_pool2d", "float32", (2, 2), (2, 2), (1, 1), (1, 1), False, True, (15, 15, 16), (0, 1),),
- ],
-)
-# fmt: on
-def test_codegen_pooling(
- typef,
- dtype,
- size,
- stride,
- dilation,
- pad,
- ceil_mode,
- count_include_pad,
- input_shape,
- expected_ops,
-):
- if skip_codegen_test():
- return
-
- low, high, _, _ = _get_low_high_atol_rtol(dtype)
- tvm_ops, acl_partitions = expected_ops
-
- shape = (1, *input_shape)
- inputs = {"a"}
- args = (shape, dtype, typef, size, stride, dilation, pad, False, False)
- func = _get_pooling_model(*args, iter(inputs))
- exp_codegen = _get_expected_pooling_codegen(*args)
-
- verify_codegen(func, exp_codegen, acl_partitions, tvm_ops)
-
-
-@pytest.mark.parametrize(
- "typef,dtype,input_shape",
- [
- ("nn.global_max_pool2d", "float32", (8, 8, 16)),
- ("nn.global_max_pool2d", "float32", (9, 9, 16)),
- ("nn.global_max_pool2d", "uint8", (8, 8, 16)),
- ("nn.global_max_pool2d", "uint8", (9, 9, 16)),
- ("nn.global_max_pool2d", "int8", (8, 8, 16)),
- ("nn.global_max_pool2d", "int8", (9, 9, 16)),
- ("nn.global_avg_pool2d", "float32", (8, 8, 16)),
- ("nn.global_avg_pool2d", "float32", (9, 9, 16)),
- ("nn.global_avg_pool2d", "uint8", (8, 8, 16)),
- ("nn.global_avg_pool2d", "uint8", (9, 9, 16)),
- ("nn.global_avg_pool2d", "int8", (8, 8, 16)),
- ("nn.global_avg_pool2d", "int8", (9, 9, 16)),
- ],
-)
-def test_codegen_global_pooling(typef, dtype, input_shape):
- if skip_codegen_test():
- return
-
- low, high, _, _ = _get_low_high_atol_rtol(dtype)
-
- shape = (1, *input_shape)
- inputs = {"a"}
- args = (shape, dtype, typef)
- func = _get_global_pooling_model(*args, iter(inputs))
- exp_codegen = _get_expected_global_pooling_codegen(*args)
- verify_codegen(func, exp_codegen, 1)
-
-
-if __name__ == "__main__":
- test_pooling()
- test_global_pooling()
- test_codegen_pooling()
- test_codegen_global_pooling()
diff --git a/tests/python/contrib/test_arm_compute_lib/test_reshape.py b/tests/python/contrib/test_arm_compute_lib/test_reshape.py
deleted file mode 100644
index 611599154c..0000000000
--- a/tests/python/contrib/test_arm_compute_lib/test_reshape.py
+++ /dev/null
@@ -1,111 +0,0 @@
-# Licensed to the Apache Software Foundation (ASF) under one
-# or more contributor license agreements. See the NOTICE file
-# distributed with this work for additional information
-# regarding copyright ownership. The ASF licenses this file
-# to you under the Apache License, Version 2.0 (the
-# "License"); you may not use this file except in compliance
-# with the License. You may obtain a copy of the License at
-#
-# http://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing,
-# software distributed under the License is distributed on an
-# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
-# KIND, either express or implied. See the License for the
-# specific language governing permissions and limitations
-# under the License.
-"""Arm Compute Library integration reshape tests."""
-
-import numpy as np
-
-import tvm
-from tvm import relay
-from tvm import testing
-
-from .infrastructure import (
- skip_runtime_test,
- skip_codegen_test,
- build_and_run,
- verify,
- verify_codegen,
-)
-from .infrastructure import Device
-
-
-def _get_model(input_shape, output_shape, dtype, var_names):
- """Return a model and any parameters it may have."""
- a = relay.var(next(var_names), shape=input_shape, dtype=dtype)
- reshape = relay.reshape(a, output_shape)
- return reshape
-
-
-def _get_expected_codegen(input_shape, output_shape, dtype):
- node = {
- "op": "kernel",
- "name": "reshape",
- "inputs": [[0, 0, 0]],
- "attrs": {
- "num_inputs": "1",
- "num_outputs": "1",
- "newshape": [[str(s) for s in output_shape]],
- "shape": [[list(output_shape)]],
- "dtype": [[dtype]],
- "allowzero": [["0"]],
- },
- }
-
- input = {
- "op": "input",
- "name": "",
- "attrs": {"shape": [[list(input_shape)]], "dtype": [[dtype]]},
- }
-
- return [input, node]
-
-
-def test_reshape():
- Device.load("test_config.json")
-
- if skip_runtime_test():
- return
-
- device = Device()
- np.random.seed(0)
-
- for dtype, low, high, atol, rtol in [
- ("float32", -127, 128, 0.001, 0.001),
- ("uint8", 0, 255, 0, 0),
- ]:
- inputs = {"a": tvm.nd.array(np.random.uniform(low, high, (1, 1, 1, 1000)).astype(dtype))}
-
- for new_shape in [(1, 1000), (10, 10, 10), (10, 100, 1), (1, 1000, 1)]:
- outputs = []
- func = _get_model(inputs["a"].shape, new_shape, dtype, iter(inputs))
- for acl in [False, True]:
- outputs.append(build_and_run(func, inputs, 1, None, device, enable_acl=acl)[0])
-
- config = {
- "new shape": inputs["a"].shape,
- "shape": new_shape,
- "dtype": dtype,
- }
- verify(outputs, atol=1e-7, rtol=1e-7, config=config)
-
-
-def test_codegen_reshape():
- if skip_codegen_test():
- return
-
- shape = (1, 1, 1, 1000)
- inputs = {"a"}
- for dtype in ["float32", "uint8"]:
- for new_shape in [(1, 1000), (10, 10, 10), (10, 100, 1)]:
- args = (shape, new_shape, dtype)
- func = _get_model(*args, iter(inputs))
- exp_codegen = _get_expected_codegen(*args)
- verify_codegen(func, exp_codegen, 1)
-
-
-if __name__ == "__main__":
- test_reshape()
- test_codegen_reshape()
diff --git a/tests/python/contrib/test_arm_compute_lib/test_runtime.py b/tests/python/contrib/test_arm_compute_lib/test_runtime.py
deleted file mode 100644
index 316dfad9e7..0000000000
--- a/tests/python/contrib/test_arm_compute_lib/test_runtime.py
+++ /dev/null
@@ -1,134 +0,0 @@
-# Licensed to the Apache Software Foundation (ASF) under one
-# or more contributor license agreements. See the NOTICE file
-# distributed with this work for additional information
-# regarding copyright ownership. The ASF licenses this file
-# to you under the Apache License, Version 2.0 (the
-# "License"); you may not use this file except in compliance
-# with the License. You may obtain a copy of the License at
-#
-# http://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing,
-# software distributed under the License is distributed on an
-# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
-# KIND, either express or implied. See the License for the
-# specific language governing permissions and limitations
-# under the License.
-"""Arm Compute Library runtime tests."""
-
-import numpy as np
-
-import tvm
-from tvm import relay
-
-from .infrastructure import skip_runtime_test, build_and_run, verify
-from .infrastructure import Device
-
-
-def test_multiple_ops():
- """
- Test multiple operators destined for ACL.
- The ACL runtime will expect these ops as 2 separate functions for
- the time being.
- """
- Device.load("test_config.json")
-
- if skip_runtime_test():
- return
-
- device = Device()
- np.random.seed(0)
-
- def get_model(input_shape, var_names):
- """Return a model and any parameters it may have."""
- a = relay.var(next(var_names), shape=input_shape, dtype="float32")
- out = relay.reshape(a, (1, 1, 1000))
- out = relay.reshape(out, (1, 1000))
- return out
-
- inputs = {"a": tvm.nd.array(np.random.uniform(0, 1, (1, 1, 1, 1000)).astype("float32"))}
-
- outputs = []
- for acl in [False, True]:
- func = get_model(inputs["a"].shape, iter(inputs))
- outputs.append(
- build_and_run(func, inputs, 1, None, device, enable_acl=acl, acl_partitions=2)[0]
- )
- verify(outputs, atol=0.002, rtol=0.01)
-
-
-def test_heterogeneous():
- """
- Test to check if offloading only supported operators works,
- while leaving unsupported operators computed via tvm.
- """
- Device.load("test_config.json")
-
- if skip_runtime_test():
- return
-
- device = Device()
- np.random.seed(0)
-
- def get_model(input_shape, var_names):
- """Return a model and any parameters it may have."""
- a = relay.var(next(var_names), shape=input_shape, dtype="float32")
- out = relay.reshape(a, (1, 1, 1000))
- out = relay.sigmoid(out)
- out = relay.reshape(out, (1, 1000))
- return out
-
- inputs = {"a": tvm.nd.array(np.random.uniform(-127, 128, (1, 1, 1, 1000)).astype("float32"))}
-
- outputs = []
- for acl in [False, True]:
- func = get_model(inputs["a"].shape, iter(inputs))
- outputs.append(
- build_and_run(
- func, inputs, 1, None, device, enable_acl=acl, tvm_ops=1, acl_partitions=2
- )[0]
- )
- verify(outputs, atol=0.002, rtol=0.01)
-
-
-def test_multiple_runs():
- """
- Test that multiple runs of an operator work.
- """
- Device.load("test_config.json")
-
- if skip_runtime_test():
- return
-
- device = Device()
-
- def get_model():
- a = relay.var("a", shape=(1, 28, 28, 512), dtype="float32")
- w = tvm.nd.array(np.ones((256, 1, 1, 512), dtype="float32"))
- weights = relay.const(w, "float32")
- conv = relay.nn.conv2d(
- a,
- weights,
- kernel_size=(1, 1),
- data_layout="NHWC",
- kernel_layout="OHWI",
- strides=(1, 1),
- padding=(0, 0),
- dilation=(1, 1),
- )
- params = {"w": w}
- return conv, params
-
- inputs = {
- "a": tvm.nd.array(np.random.uniform(-127, 128, (1, 28, 28, 512)).astype("float32")),
- }
-
- func, params = get_model()
- outputs = build_and_run(func, inputs, 1, params, device, enable_acl=True, no_runs=3)
- verify(outputs, atol=0.002, rtol=0.01)
-
-
-if __name__ == "__main__":
- test_multiple_ops()
- test_heterogeneous()
- test_multiple_runs()
diff --git a/tests/python/runtime/test_runtime_rpc.py b/tests/python/runtime/test_runtime_rpc.py
index 759c7a2e5c..a439be9ef8 100644
--- a/tests/python/runtime/test_runtime_rpc.py
+++ b/tests/python/runtime/test_runtime_rpc.py
@@ -250,7 +250,9 @@ def test_rpc_remote_module():
n = tvm.runtime.convert(102)
A = te.placeholder((n,), name="A")
B = te.compute(A.shape, lambda *i: A(*i) + 1.0, name="B")
- mod = tvm.ir.IRModule({"main": te.create_prim_func([A, B])})
+ mod = tvm.ir.IRModule.from_expr(
+ te.create_prim_func([A, B]).with_attr("global_symbol", "myadd")
+ )
server0 = rpc.Server(key="x0")
server1 = rpc.Server(key="x1")
@@ -265,7 +267,7 @@ def test_rpc_remote_module():
def check_remote(remote):
temp = utils.tempdir()
dev = remote.cpu(0)
- f = tvm.build(mod, [A, B], "llvm", name="myadd")
+ f = tvm.build(mod, "llvm")
path_dso = temp.relpath("dev_lib.so")
f.export_library(path_dso)
remote.upload(path_dso)
@@ -296,7 +298,7 @@ def test_rpc_remote_module():
# export to minrpc
temp = utils.tempdir()
# system lib prefix will trigger system lib build
- f = tvm.build(mod.with_attr("system_lib_prefix", ""), [A, B], "llvm", name="myadd")
+ f = tvm.build(mod.with_attr("system_lib_prefix", ""), "llvm")
path_minrpc = temp.relpath("dev_lib.minrpc")
f.export_library(path_minrpc, fcompile=rpc.with_minrpc(cc.create_executable))
@@ -339,25 +341,7 @@ def test_rpc_remote_module():
xo, xi = s.split(x, factors=[None ,32])
s.bind(xo, "blockIdx.x")
s.bind(xi, "threadIdx.x")
- f = tvm.build(s, [A, B], "opencl --host=llvm", name="myadd")
- # Option 1: save modules separately and rely on remote compiler
- path_o = temp.relpath("myadd.o")
- path_cl = temp.relpath("myadd.cl")
- path_json = temp.relpath("myadd.tvm_meta.json")
- f.save(path_o)
- f.imported_modules[0].save(path_cl)
- remote.upload(path_o)
- remote.upload(path_cl)
- # upload meta data
- remote.upload(path_json)
- fhost = remote.load_module("myadd.o")
- fdev = remote.load_module("myadd.cl")
- fhost.import_module(fdev)
- a = tvm.nd.array(np.random.uniform(size=102).astype(A.dtype), dev)
- b = tvm.nd.array(np.zeros(102, dtype=A.dtype), dev)
- fhost(a, b)
- np.testing.assert_equal(b.numpy(), a.numpy() + 1)
- # Option 2: export library as a tar ball then handled by remote compiler
+ f = tvm.build(s.mod, "opencl --host=llvm")
path_tar = temp.relpath("myadd.tar")
f.export_library(path_tar)
remote.upload(path_tar)
diff --git a/tests/scripts/ci.py b/tests/scripts/ci.py
index 344c103850..0bd97e4ee0 100755
--- a/tests/scripts/ci.py
+++ b/tests/scripts/ci.py
@@ -594,7 +594,6 @@ generated = [
help="Run GPU build and test(s)",
options={
"cpp": CPP_UNITTEST,
- "topi": ("run topi tests", ["./tests/scripts/task_python_topi.sh"]),
"unittest": (
"run unit tests",
[
diff --git a/tests/scripts/task_java_unittest.sh b/tests/scripts/task_java_unittest.sh
index ef13e79638..a35b023ad0 100755
--- a/tests/scripts/task_java_unittest.sh
+++ b/tests/scripts/task_java_unittest.sh
@@ -37,7 +37,6 @@ trap cleanup 0
python3 "$SCRIPT_DIR"/test_add_cpu.py "$TEMP_DIR"
python3 "$SCRIPT_DIR"/test_add_gpu.py "$TEMP_DIR"
-python3 "$SCRIPT_DIR"/test_graph_executor.py "$TEMP_DIR"
# Skip the Java RPC Unittests, see https://github.com/apache/tvm/issues/13168
# # start rpc proxy server
diff --git a/tests/scripts/task_python_arm_compute_library.sh b/tests/scripts/task_python_arm_compute_library.sh
index 520abb58cd..b00f1ee423 100755
--- a/tests/scripts/task_python_arm_compute_library.sh
+++ b/tests/scripts/task_python_arm_compute_library.sh
@@ -26,4 +26,3 @@ find . -type f -path "*.pyc" | xargs rm -f
# setup cython
cd python; python3 setup.py build_ext --inplace; cd ..
-run_pytest python-arm_compute_lib tests/python/contrib/test_arm_compute_lib
diff --git a/tests/scripts/task_python_integration.sh b/tests/scripts/task_python_integration.sh
index ecab6b7267..326743394d 100755
--- a/tests/scripts/task_python_integration.sh
+++ b/tests/scripts/task_python_integration.sh
@@ -35,12 +35,3 @@ find . -type f -path "*.pyc" | xargs rm -f
# setup cython
cd python; python3 setup.py build_ext --inplace; cd ..
-
-run_pytest ${TVM_INTEGRATION_TESTSUITE_NAME}-integration tests/python/integration
-
-
-# Command line driver test
-run_pytest ${TVM_INTEGRATION_TESTSUITE_NAME}-driver tests/python/driver
-
-# Target test
-run_pytest ${TVM_INTEGRATION_TESTSUITE_NAME}-target tests/python/target
diff --git a/tests/scripts/task_python_topi.sh b/tests/scripts/task_python_topi.sh
index 97030d6948..9cb1fd476e 100755
--- a/tests/scripts/task_python_topi.sh
+++ b/tests/scripts/task_python_topi.sh
@@ -25,11 +25,3 @@ export TVM_TEST_TARGETS="llvm; cuda"
# to avoid CI thread throttling.
export TVM_BIND_THREADS=0
export OMP_NUM_THREADS=1
-
-# setup cython
-cd python; python3 setup.py build_ext --inplace; cd ..
-
-# cleanup pycache
-find . -type f -path "*.pyc" | xargs rm -f
-
-run_pytest python-topi tests/python/topi/
diff --git a/tests/scripts/task_python_unittest.sh b/tests/scripts/task_python_unittest.sh
index b84694cfd5..b074af3147 100755
--- a/tests/scripts/task_python_unittest.sh
+++ b/tests/scripts/task_python_unittest.sh
@@ -37,12 +37,9 @@ run_pytest ${TVM_UNITTEST_TESTSUITE_NAME}-platform-minimal-test tests/python/all
# Then run all unittests on both ctypes and cython.
TEST_FILES=(
"arith"
- "auto_scheduler"
- "autotvm"
"codegen"
"ir"
"meta_schedule"
- "micro"
"runtime"
"te"
"testing"
@@ -54,6 +51,7 @@ TEST_FILES=(
"tvmscript"
"usmp"
"ci"
+ "target"
)
for TEST_FILE in ${TEST_FILES[@]}; do