From f0cf278f81ac5f8b117651064f4c73f160cd18c7 Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Tue, 4 Oct 2022 09:48:48 -0700 Subject: [PATCH 1/9] remove target from test functions --- tests/python/contrib/test_hexagon/test_benchmark_maxpool2d.py | 1 - tests/python/contrib/test_hexagon/test_sigmoid.py | 1 - tests/python/contrib/test_hexagon/topi/test_conv2d_slice.py | 1 - tests/python/contrib/test_hexagon/topi/test_dwconv2d_slice.py | 1 - tests/python/contrib/test_hexagon/topi/test_relu_slice.py | 1 - 5 files changed, 5 deletions(-) diff --git a/tests/python/contrib/test_hexagon/test_benchmark_maxpool2d.py b/tests/python/contrib/test_hexagon/test_benchmark_maxpool2d.py index 41169494417a..26352d4ae51b 100644 --- a/tests/python/contrib/test_hexagon/test_benchmark_maxpool2d.py +++ b/tests/python/contrib/test_hexagon/test_benchmark_maxpool2d.py @@ -247,7 +247,6 @@ def test_maxpool2d_nhwc( ) target_hexagon = tvm.target.hexagon("v69", link_params=True) - # func = tvm.build(sch.mod, target=tvm.target.Target(target_hexagon, host=target_hexagon)) built_module = tvm.build( sch.mod, target=tvm.target.Target(target_hexagon, host=target_hexagon) ) diff --git a/tests/python/contrib/test_hexagon/test_sigmoid.py b/tests/python/contrib/test_hexagon/test_sigmoid.py index 9aad35ee76c1..1f9f5501d816 100644 --- a/tests/python/contrib/test_hexagon/test_sigmoid.py +++ b/tests/python/contrib/test_hexagon/test_sigmoid.py @@ -71,7 +71,6 @@ def test_sigmoid( dtype, input_np, ref_output_np, - target, hexagon_session, ): InputTensor = te.placeholder(in_shape, name="InputTensor", dtype=dtype) diff --git a/tests/python/contrib/test_hexagon/topi/test_conv2d_slice.py b/tests/python/contrib/test_hexagon/topi/test_conv2d_slice.py index a03c35cb9e78..b3ef6890e290 100755 --- a/tests/python/contrib/test_hexagon/topi/test_conv2d_slice.py +++ b/tests/python/contrib/test_hexagon/topi/test_conv2d_slice.py @@ -274,7 +274,6 @@ def test_conv2d( input_np_padded, weights_np_transformed, expected_output_np, - target, working_scope, hexagon_session, ): diff --git a/tests/python/contrib/test_hexagon/topi/test_dwconv2d_slice.py b/tests/python/contrib/test_hexagon/topi/test_dwconv2d_slice.py index f8419e9c55d8..bc72e2bf4bbc 100644 --- a/tests/python/contrib/test_hexagon/topi/test_dwconv2d_slice.py +++ b/tests/python/contrib/test_hexagon/topi/test_dwconv2d_slice.py @@ -258,7 +258,6 @@ def test_dwconv2d( input_np_padded, weights_np_transformed, expected_output_np, - target, working_scope, hexagon_session, ): diff --git a/tests/python/contrib/test_hexagon/topi/test_relu_slice.py b/tests/python/contrib/test_hexagon/topi/test_relu_slice.py index c08d4a5545f1..87978f586541 100644 --- a/tests/python/contrib/test_hexagon/topi/test_relu_slice.py +++ b/tests/python/contrib/test_hexagon/topi/test_relu_slice.py @@ -75,7 +75,6 @@ def test_relu( output_layout, transformed_input_np, transformed_ref_output_np, - target, working_scope, hexagon_session, ): From 831e38190b18796a00368f139653542406fb67a2 Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Tue, 4 Oct 2022 10:40:28 -0700 Subject: [PATCH 2/9] refactor target_hexagon --- .../contrib/test_hexagon/infrastructure.py | 6 +++++ .../contrib/test_hexagon/test_autotvm.py | 11 +++++---- .../test_hexagon/test_benchmark_maxpool2d.py | 11 +++------ .../test_hexagon/test_cache_read_write.py | 5 ++-- .../test_hexagon/test_fixed_point_multiply.py | 8 +++---- .../contrib/test_hexagon/test_launcher.py | 17 ++++++-------- .../test_hexagon/test_meta_schedule.py | 23 ++++++------------- .../contrib/test_hexagon/test_models.py | 5 ++-- .../contrib/test_hexagon/test_parallel_hvx.py | 11 ++++----- .../test_parallel_hvx_load_vtcm.py | 16 +++++-------- .../test_hexagon/test_parallel_scalar.py | 11 ++++----- .../contrib/test_hexagon/test_sigmoid.py | 7 ++---- .../test_software_pipeline_async.py | 7 +++--- .../contrib/test_hexagon/test_thread_pool.py | 10 ++++---- .../test_hexagon/test_vtcm_bandwidth.py | 12 ++++------ .../topi/test_add_subtract_multiply.py | 10 +++++--- .../test_hexagon/topi/test_argmax_slice.py | 6 ++--- .../topi/test_avg_pool2d_slice.py | 12 ++++++---- .../test_hexagon/topi/test_batch_matmul.py | 5 ++-- .../test_hexagon/topi/test_cast_slice.py | 14 +++++------ .../contrib/test_hexagon/topi/test_clip.py | 9 +++----- .../topi/test_conv2d_fp16_intrin.py | 7 +++--- .../test_hexagon/topi/test_conv2d_nchw.py | 3 ++- .../test_hexagon/topi/test_conv2d_nhwc.py | 6 ++--- .../test_hexagon/topi/test_conv2d_slice.py | 7 ++---- .../topi/test_conv2d_transpose.py | 6 +++-- .../contrib/test_hexagon/topi/test_dense.py | 6 ++--- .../test_hexagon/topi/test_depth_to_space.py | 7 ++---- .../topi/test_depthwise_conv2d.py | 3 ++- .../topi/test_dequantize_slice.py | 11 +++++---- .../test_hexagon/topi/test_dwconv2d_slice.py | 7 ++---- .../topi/test_max_pool2d_slice.py | 9 ++------ .../contrib/test_hexagon/topi/test_pad.py | 4 +++- .../contrib/test_hexagon/topi/test_pooling.py | 6 +++-- .../test_hexagon/topi/test_quantize.py | 11 +++++---- .../contrib/test_hexagon/topi/test_reduce.py | 5 ++-- .../test_hexagon/topi/test_relu_slice.py | 8 ++----- .../contrib/test_hexagon/topi/test_reshape.py | 10 +++----- .../test_hexagon/topi/test_resize2d.py | 5 ++-- .../contrib/test_hexagon/topi/test_softmax.py | 6 ++--- .../test_hexagon/topi/test_softmax_slice.py | 3 ++- .../test_hexagon/topi/test_tanh_slice.py | 6 ++--- 42 files changed, 162 insertions(+), 190 deletions(-) diff --git a/tests/python/contrib/test_hexagon/infrastructure.py b/tests/python/contrib/test_hexagon/infrastructure.py index 1058e1dd8117..1d854b5fd64f 100644 --- a/tests/python/contrib/test_hexagon/infrastructure.py +++ b/tests/python/contrib/test_hexagon/infrastructure.py @@ -351,3 +351,9 @@ def quantize_np(arr_np: numpy.ndarray, dtype: str): zero_point = numpy.rint((fmax * qmin - fmin * qmax) / (fmax - fmin)).astype("int32") quant_np = (arr_np / scale + zero_point).astype(dtype) return quant_np, scale, zero_point + + +def get_hexagon_target(cpu_ver: str) -> tvm.target.Target: + """Creates a Hexagon target""" + target = tvm.target.hexagon(cpu_ver) + return tvm.target.Target(target, host=target) diff --git a/tests/python/contrib/test_hexagon/test_autotvm.py b/tests/python/contrib/test_hexagon/test_autotvm.py index 513d5bdbab7a..da60e20c3bf4 100644 --- a/tests/python/contrib/test_hexagon/test_autotvm.py +++ b/tests/python/contrib/test_hexagon/test_autotvm.py @@ -19,8 +19,6 @@ import contextlib import os -import sys - import pytest import tvm @@ -28,6 +26,8 @@ from tvm import autotvm, te from tvm.autotvm.tuner import GATuner, XGBTuner +from .infrastructure import get_hexagon_target + @autotvm.template("demo_template") def demo_template(): @@ -143,12 +143,13 @@ def test_autotvm(hexagon_session): ), ), } - target_hexagon = tvm.target.hexagon("v68") task = autotvm.task.create( - "demo_template", args=[], target=target_hexagon, target_host=target_hexagon + "demo_template", + args=[], + target=get_hexagon_target("v68"), ) tune_tasks([task], **options) if __name__ == "__main__": - sys.exit(pytest.main(sys.argv)) + tvm.testing.main() diff --git a/tests/python/contrib/test_hexagon/test_benchmark_maxpool2d.py b/tests/python/contrib/test_hexagon/test_benchmark_maxpool2d.py index 26352d4ae51b..b2de2c7e95f0 100644 --- a/tests/python/contrib/test_hexagon/test_benchmark_maxpool2d.py +++ b/tests/python/contrib/test_hexagon/test_benchmark_maxpool2d.py @@ -42,7 +42,6 @@ primfuncs and demonstrate more coding strategies. """ -import sys import pytest import numpy as np import copy @@ -51,15 +50,11 @@ import tvm.testing from tvm import te, topi, tir from tvm.topi import testing -from tvm.script import tir as T -from tvm.tir import IndexMap -from tvm.relay.backend import Executor, Runtime from tvm.contrib.hexagon.session import Session from typing import List -from .infrastructure import allocate_hexagon_array +from .infrastructure import allocate_hexagon_array, get_hexagon_target from . import benchmark_util as bu -from .benchmark_util import benchmark_group _SHOULD_SKIP_BENCHMARKS, _SKIP_BENCHMARKS_REASON = bu.skip_bencharks_flag_and_reason() @@ -246,9 +241,9 @@ def test_maxpool2d_nhwc( block="tensor", buffer="placeholder", index_map=_int8_nhwc_8h8w32c_map ) - target_hexagon = tvm.target.hexagon("v69", link_params=True) built_module = tvm.build( - sch.mod, target=tvm.target.Target(target_hexagon, host=target_hexagon) + sch.mod, + target=get_hexagon_target("v69"), ) # Save a local copy of the Hexagon object code (in the form of a .so file) diff --git a/tests/python/contrib/test_hexagon/test_cache_read_write.py b/tests/python/contrib/test_hexagon/test_cache_read_write.py index 896db8b59c5c..af5e7a398870 100644 --- a/tests/python/contrib/test_hexagon/test_cache_read_write.py +++ b/tests/python/contrib/test_hexagon/test_cache_read_write.py @@ -24,6 +24,8 @@ from tvm.contrib.hexagon.session import Session from tvm.script import tir as T +from .infrastructure import get_hexagon_target + def intrin_mem_copy(shape, dtype, dst_scope, src_scope): """Define and return tensor intrinsic for mem copy""" @@ -76,11 +78,10 @@ def verify(hexagon_session: Session, schedule, x_tensor, y_tensor, z_tensor, siz """Verify correctness with reference from numpy""" print(tvm.lower(schedule, [x_tensor, y_tensor, z_tensor])) - target_hexagon = tvm.target.hexagon("v68", link_params=True) func = tvm.build( schedule, [x_tensor, y_tensor, z_tensor], - tvm.target.Target(target_hexagon, host=target_hexagon), + get_hexagon_target("v68"), name="dmacpy", ) diff --git a/tests/python/contrib/test_hexagon/test_fixed_point_multiply.py b/tests/python/contrib/test_hexagon/test_fixed_point_multiply.py index 8ee04a649990..ee03599ff1f4 100644 --- a/tests/python/contrib/test_hexagon/test_fixed_point_multiply.py +++ b/tests/python/contrib/test_hexagon/test_fixed_point_multiply.py @@ -14,14 +14,15 @@ # KIND, either express or implied. See the License for the # specific language governing permissions and limitations # under the License. +import re +import numpy as np import tvm.testing from tvm import relay from tvm.relay.backend import Executor from tvm.contrib.hexagon.session import Session -import re -import numpy as np +from .infrastructure import get_hexagon_target @tvm.testing.requires_hexagon @@ -37,13 +38,12 @@ def test_vmpy_intrinsic_presence(): relay_mod = tvm.IRModule.from_expr(y) params = {} - target_hexagon = tvm.target.hexagon("v68") executor = Executor("graph", {"link-params": True}) with tvm.transform.PassContext(opt_level=3): hexagon_lowered = tvm.relay.build( relay_mod, - tvm.target.Target(target_hexagon, host=target_hexagon), + get_hexagon_target("v68"), executor=executor, params=params, ) diff --git a/tests/python/contrib/test_hexagon/test_launcher.py b/tests/python/contrib/test_hexagon/test_launcher.py index 9321ddf71d3b..f606b9082720 100644 --- a/tests/python/contrib/test_hexagon/test_launcher.py +++ b/tests/python/contrib/test_hexagon/test_launcher.py @@ -24,6 +24,8 @@ from tvm.contrib.hexagon.session import Session from tvm.relay.backend import Executor, Runtime +from .infrastructure import get_hexagon_target + @tvm.testing.requires_hexagon def test_add(hexagon_session: Session): @@ -36,11 +38,10 @@ def test_add(hexagon_session: Session): ) sched = tvm.te.create_schedule(compute_c.op) - target_hexagon = tvm.target.hexagon("v68", link_params=True) func = tvm.build( sched, [placeholder_a, placeholder_b, compute_c], - tvm.target.Target(target_hexagon, host=target_hexagon), + get_hexagon_target("v68"), name="add", ) @@ -67,11 +68,10 @@ def test_add_vtcm(hexagon_session: Session): ) sched = tvm.te.create_schedule(compute_c.op) - target_hexagon = tvm.target.hexagon("v68", link_params=True) func = tvm.build( sched, [placeholder_a, placeholder_b, compute_c], - tvm.target.Target(target_hexagon, host=target_hexagon), + get_hexagon_target("v68"), name="add", ) @@ -116,11 +116,10 @@ def test_matmul(self, hexagon_session, size_m, size_n, size_k): ) schedule = te.create_schedule(compute_z.op) - target_hexagon = tvm.target.hexagon("v68", link_params=True) func = tvm.build( schedule, [placeholder_x, placeholder_y, compute_z], - tvm.target.Target(target_hexagon, host=target_hexagon), + get_hexagon_target("v68"), ) mod = hexagon_session.load_module(func) @@ -172,7 +171,6 @@ def test_graph_executor(hexagon_session: Session): relay_mod = tvm.IRModule.from_expr(f) relay_mod = relay.transform.InferType()(relay_mod) - target_hexagon = tvm.target.hexagon("v68") runtime = Runtime("cpp") executor = Executor("graph") @@ -184,7 +182,7 @@ def test_graph_executor(hexagon_session: Session): with tvm.transform.PassContext(opt_level=3): lowered = tvm.relay.build( relay_mod, - tvm.target.Target(target_hexagon, host=target_hexagon), + get_hexagon_target("v68"), runtime=runtime, executor=executor, ) @@ -242,14 +240,13 @@ def test_graph_executor_multiple_conv2d(hexagon_session: Session): relay_mod = tvm.IRModule.from_expr(f) relay_mod = relay.transform.InferType()(relay_mod) - target_hexagon = tvm.target.hexagon("v68") runtime = Runtime("cpp") executor = Executor("graph") with tvm.transform.PassContext(opt_level=3): lowered = tvm.relay.build( relay_mod, - tvm.target.Target(target_hexagon, host=target_hexagon), + get_hexagon_target("v68"), runtime=runtime, executor=executor, ) diff --git a/tests/python/contrib/test_hexagon/test_meta_schedule.py b/tests/python/contrib/test_hexagon/test_meta_schedule.py index 74f3ab673ec8..8b07122c2a17 100644 --- a/tests/python/contrib/test_hexagon/test_meta_schedule.py +++ b/tests/python/contrib/test_hexagon/test_meta_schedule.py @@ -33,8 +33,8 @@ from tvm.meta_schedule.runner import RunnerInput from tvm.contrib.hexagon.meta_schedule import get_hexagon_local_builder, get_hexagon_rpc_runner from tvm.relay.backend import Executor -from tvm.topi.utils import get_const_tuple -from tvm.meta_schedule.testing import te_workload + +from .infrastructure import get_hexagon_target MATMUL_N = 16 MATMUL_M = 32 @@ -61,14 +61,12 @@ def test_builder_runner(hexagon_launcher): if hexagon_launcher._serial_number == "simulator": pytest.skip(msg="Tuning on simulator not supported.") - target_hexagon = tvm.target.hexagon("v68", link_params=True) - target = tvm.target.Target(target_hexagon, host=target_hexagon) mod = MatmulModule builder = get_hexagon_local_builder() runner = get_hexagon_rpc_runner(hexagon_launcher, number=1, repeat=1, min_repeat_ms=0) - (builder_result,) = builder.build([BuilderInput(mod, target)]) + (builder_result,) = builder.build([BuilderInput(mod, get_hexagon_target("v68"))]) assert builder_result.artifact_path is not None assert builder_result.error_msg is None @@ -177,8 +175,6 @@ def test_vrmpy_dense(hexagon_launcher): pytest.skip(msg="Tuning on simulator not supported.") do_tune = True - target_hexagon = tvm.target.hexagon("v68") - target = tvm.target.Target(target_hexagon, host=target_hexagon) M, N, K = 128, 768, 768 workload = te.create_prim_func(dense(M, N, K)) @@ -212,7 +208,7 @@ def schedule_dense_for_tune(sch): ) with hexagon_launcher.start_session() as session: - verify_dense(sch, target, M, N, K, session) + verify_dense(sch, get_hexagon_target("v68"), M, N, K, session) # This is an example of a schedule found by vrmpy auto tensorization. @@ -274,9 +270,6 @@ def test_vrmpy_dense_auto_tensorize(hexagon_launcher): if hexagon_launcher._serial_number == "simulator": pytest.skip(msg="Tuning on simulator not supported.") - target_hexagon = tvm.target.hexagon("v68") - target = tvm.target.Target(target_hexagon, host=target_hexagon) - M, N, K = 128, 768, 768 workload = te.create_prim_func(dense(M, N, K)) @@ -319,7 +312,7 @@ def test_vrmpy_dense_auto_tensorize(hexagon_launcher): sch = ms.tune_tir( mod=workload, - target=target, + target=get_hexagon_target("v68"), config=config, work_dir=work_dir, sch_rules=lambda: sch_rules, @@ -331,7 +324,7 @@ def test_vrmpy_dense_auto_tensorize(hexagon_launcher): sch = tvm.tir.Schedule(Module_vrmpy_auto_tensorize, debug_mask="all") with hexagon_launcher.start_session() as session: - verify_dense(sch, target, M, N, K, session) + verify_dense(sch, get_hexagon_target("v68"), M, N, K, session) @tvm.testing.requires_hexagon @@ -339,8 +332,6 @@ def test_conv2d_relay_auto_schedule(hexagon_launcher): if hexagon_launcher._serial_number == "simulator": pytest.skip(msg="Tuning on simulator not supported.") - target_hexagon = tvm.target.hexagon("v69") - target = tvm.target.Target(target_hexagon, host=target_hexagon) I, O, H, W = 64, 64, 56, 56 kH = kW = 3 @@ -400,7 +391,7 @@ def test_conv2d_relay_auto_schedule(hexagon_launcher): lib = ms.tune_relay( mod=mod, params=params, - target=target, + target=get_hexagon_target("v69"), config=config, work_dir=work_dir, builder=get_hexagon_local_builder(), diff --git a/tests/python/contrib/test_hexagon/test_models.py b/tests/python/contrib/test_hexagon/test_models.py index 78d7116d0853..95e5191a8619 100644 --- a/tests/python/contrib/test_hexagon/test_models.py +++ b/tests/python/contrib/test_hexagon/test_models.py @@ -25,6 +25,8 @@ from tvm.contrib.hexagon.session import Session from tvm.relay.backend import Executor, Runtime +from .infrastructure import get_hexagon_target + def get_mobilenet(): """Download and import mobilenet model with ONNX""" @@ -43,7 +45,6 @@ def test_mobilenet(hexagon_session: Session): dtype = "float32" onnx_model = get_mobilenet() - target_hexagon = tvm.target.hexagon("v68") target_llvm = tvm.target.Target("llvm") runtime = Runtime("cpp") executor = Executor("graph", {"link-params": True}) @@ -58,7 +59,7 @@ def test_mobilenet(hexagon_session: Session): with tvm.transform.PassContext(opt_level=3): hexagon_lowered = tvm.relay.build( relay_mod, - tvm.target.Target(target_hexagon, host=target_hexagon), + get_hexagon_target("v68"), runtime=runtime, executor=executor, params=params, diff --git a/tests/python/contrib/test_hexagon/test_parallel_hvx.py b/tests/python/contrib/test_hexagon/test_parallel_hvx.py index a34f5b8e261b..75ba885c8fff 100644 --- a/tests/python/contrib/test_hexagon/test_parallel_hvx.py +++ b/tests/python/contrib/test_hexagon/test_parallel_hvx.py @@ -19,10 +19,12 @@ Test parallelizing HVX workloads and compare them to single thread examples. """ import numpy as np -import tvm +from numpy.random import default_rng +import tvm from tvm.script import tir as T -from numpy.random import default_rng + +from .infrastructure import get_hexagon_target TEST_OUTPUT_TEMPLATE = "Test {} with {} operations... \n -Single Thread: {} ms \n -Parallel: {} ms\n -Speedup: {}x\n" @@ -132,10 +134,7 @@ def operator(a: T.handle, b: T.handle, c: T.handle) -> None: def evaluate(hexagon_session, shape_dtypes, expected_output_producer, sch): a_shape, a_dtype, b_shape, b_dtype, c_shape, c_dtype = shape_dtypes - target_hexagon = tvm.target.hexagon("v68") - func_tir = tvm.build( - sch.mod["main"], target=tvm.target.Target(target_hexagon, host=target_hexagon) - ) + func_tir = tvm.build(sch.mod["main"], get_hexagon_target("v68")) module = hexagon_session.load_module(func_tir) rng = default_rng() diff --git a/tests/python/contrib/test_hexagon/test_parallel_hvx_load_vtcm.py b/tests/python/contrib/test_hexagon/test_parallel_hvx_load_vtcm.py index 5dcb4b18b845..6e43298a4eb5 100644 --- a/tests/python/contrib/test_hexagon/test_parallel_hvx_load_vtcm.py +++ b/tests/python/contrib/test_hexagon/test_parallel_hvx_load_vtcm.py @@ -18,10 +18,12 @@ """ Test different strategies for loading data into vtcm before running HVX workloads. """ import numpy as np -import tvm +from numpy.random import default_rng +import tvm from tvm.script import tir as T -from numpy.random import default_rng + +from .infrastructure import get_hexagon_target TEST_OUTPUT_TEMPLATE = "Test with {} MB of data to load... \n -No VTCM: {} Gops \n -Basic VTCM: {} Gops \n -Vectorized: {} Gops\n -Vectorized and Parallelized: {} Gops\n -Preallocated and Vectorized: {} Gops\n -Preallocated, Vectorized, and Parallelized: {} Gops\n -Single DMA: {} Gops\n -Preloaded: {} Gops\n" @@ -299,10 +301,7 @@ def evaluate_result(operations, tag, time, result, expected_output): def setup_and_run(hexagon_session, sch, a, b, c, operations, mem_scope="global"): - target_hexagon = tvm.target.hexagon("v69") - func_tir = tvm.build( - sch.mod["main"], target=tvm.target.Target(target_hexagon, host=target_hexagon) - ) + func_tir = tvm.build(sch.mod["main"], target=get_hexagon_target("v69")) module = hexagon_session.load_module(func_tir) a_hexagon = tvm.runtime.ndarray.array(a, device=hexagon_session.device, mem_scope=mem_scope) @@ -322,10 +321,7 @@ def setup_and_run(hexagon_session, sch, a, b, c, operations, mem_scope="global") def setup_and_run_preallocated(hexagon_session, sch, a, b, c, operations): - target_hexagon = tvm.target.hexagon("v69") - func_tir = tvm.build( - sch.mod["main"], target=tvm.target.Target(target_hexagon, host=target_hexagon) - ) + func_tir = tvm.build(sch.mod["main"], target=get_hexagon_target("v69")) module = hexagon_session.load_module(func_tir) a_vtcm = np.zeros((a.size), dtype="uint8") diff --git a/tests/python/contrib/test_hexagon/test_parallel_scalar.py b/tests/python/contrib/test_hexagon/test_parallel_scalar.py index b3d07ae978ba..fd3eef1b195b 100644 --- a/tests/python/contrib/test_hexagon/test_parallel_scalar.py +++ b/tests/python/contrib/test_hexagon/test_parallel_scalar.py @@ -18,10 +18,12 @@ """ Test parallelism for multiple different scalar workloads. """ import numpy as np -import tvm +from numpy.random import default_rng +import tvm from tvm.script import tir as T -from numpy.random import default_rng + +from .infrastructure import get_hexagon_target TEST_OUTPUT_TEMPLATE = "Test {} with {} operations... \n -Single Thread: {} ms \n -Parallel: {} ms\n -Speedup: {}x\n" @@ -75,10 +77,7 @@ def evaluate(hexagon_session, operations, expected, sch): shape = operations dtype = "float64" - target_hexagon = tvm.target.hexagon("v68") - func_tir = tvm.build( - sch.mod["main"], target=tvm.target.Target(target_hexagon, host=target_hexagon) - ) + func_tir = tvm.build(sch.mod["main"], target=get_hexagon_target("v68")) module = hexagon_session.load_module(func_tir) rng = default_rng() diff --git a/tests/python/contrib/test_hexagon/test_sigmoid.py b/tests/python/contrib/test_hexagon/test_sigmoid.py index 1f9f5501d816..1ff5bf3db340 100644 --- a/tests/python/contrib/test_hexagon/test_sigmoid.py +++ b/tests/python/contrib/test_hexagon/test_sigmoid.py @@ -25,7 +25,7 @@ from tvm import topi from tvm.contrib.hexagon.build import HexagonLauncher -from .infrastructure import allocate_hexagon_array, transform_numpy +from .infrastructure import allocate_hexagon_array, get_hexagon_target def sigmoid_compute(Input): @@ -77,9 +77,6 @@ def test_sigmoid( OutputTensor = sigmoid_compute(InputTensor) - target_hexagon = tvm.target.hexagon("v69") - target = tvm.target.Target(target_hexagon, host=target_hexagon) - tir_s = sigmoid_stir_schedule(InputTensor, OutputTensor) input_data = allocate_hexagon_array( @@ -94,7 +91,7 @@ def test_sigmoid( func_name = "sigmoid" with tvm.transform.PassContext(opt_level=3): - runtime_module = tvm.build(tir_s.mod, target=target, name=func_name) + runtime_module = tvm.build(tir_s.mod, target=get_hexagon_target("v69"), name=func_name) assert "hvx_sigmoid" in runtime_module.get_source("asm") assert "vmin" in runtime_module.get_source("asm") diff --git a/tests/python/contrib/test_hexagon/test_software_pipeline_async.py b/tests/python/contrib/test_hexagon/test_software_pipeline_async.py index 6bcca90ec9d3..5ae67e6501e0 100644 --- a/tests/python/contrib/test_hexagon/test_software_pipeline_async.py +++ b/tests/python/contrib/test_hexagon/test_software_pipeline_async.py @@ -24,6 +24,8 @@ from tvm.contrib.hexagon.session import Session from tvm.script import tir as T +from .infrastructure import get_hexagon_target + outer = tvm.testing.parameter(8, 16) inner = tvm.testing.parameter(64, 128) scope = tvm.testing.parameter("global", "global.vtcm") @@ -63,11 +65,8 @@ def test_software_pipeline_with_cache_read(hexagon_launcher, compute, outer, inn b_np = np.random.uniform(low=0, high=128, size=(outer, inner)).astype(dtype) ref = compute[1](a_np) - target_hexagon = tvm.target.hexagon("v68", link_params=True) with tvm.transform.PassContext(config={"tir.use_async_copy": 1}): - func = tvm.build( - sch.mod["main"], target=tvm.target.Target(target_hexagon, host=target_hexagon) - ) + func = tvm.build(sch.mod["main"], target=get_hexagon_target("v68")) with hexagon_launcher.start_session() as hexagon_session: dev = hexagon_session.device diff --git a/tests/python/contrib/test_hexagon/test_thread_pool.py b/tests/python/contrib/test_hexagon/test_thread_pool.py index c943dbbb22e5..2fc82cf49984 100644 --- a/tests/python/contrib/test_hexagon/test_thread_pool.py +++ b/tests/python/contrib/test_hexagon/test_thread_pool.py @@ -26,6 +26,8 @@ from tvm.contrib.hexagon.session import Session from tvm.script import tir as T +from .infrastructure import get_hexagon_target + @tvm.script.ir_module class ElemwiseSumIRModule: @@ -73,9 +75,9 @@ def benchmark_func(mod, name, args, hexagon_session): @tvm.testing.requires_hexagon def test_speedup(hexagon_session: Session, capsys): """Test speedup""" - target_hexagon = tvm.target.hexagon("v68", link_params=True) func = tvm.build( - ElemwiseSumIRModule, target=tvm.target.Target(target_hexagon, host=target_hexagon) + ElemwiseSumIRModule, + target=get_hexagon_target("v68"), ) mod = hexagon_session.load_module(func) args = generate_add_test_data(hexagon_session) @@ -89,9 +91,9 @@ def test_speedup(hexagon_session: Session, capsys): @tvm.testing.requires_hexagon def test_elemwise_sum_parallel(hexagon_session: Session): """Test parallel elementwise sum""" - target_hexagon = tvm.target.hexagon("v68", link_params=True) func = tvm.build( - ElemwiseSumIRModule, target=tvm.target.Target(target_hexagon, host=target_hexagon) + ElemwiseSumIRModule, + target=get_hexagon_target("v68"), ) mod = hexagon_session.load_module(func) diff --git a/tests/python/contrib/test_hexagon/test_vtcm_bandwidth.py b/tests/python/contrib/test_hexagon/test_vtcm_bandwidth.py index 83daf2458737..307d3a96bf15 100644 --- a/tests/python/contrib/test_hexagon/test_vtcm_bandwidth.py +++ b/tests/python/contrib/test_hexagon/test_vtcm_bandwidth.py @@ -18,11 +18,12 @@ """Test theoretical bandwith for data transfers to VTCM for different strategies.""" import numpy as np -from tests.python.contrib.test_hexagon.infrastructure import allocate_hexagon_array -import tvm +from numpy.random import default_rng +import tvm from tvm.script import tir as T -from numpy.random import default_rng + +from .infrastructure import get_hexagon_target MB = 1024**2 KB = 1024 @@ -81,10 +82,7 @@ def operator(a: T.handle, a_v: T.handle) -> None: def evaluate(hexagon_session, sch, size): a_shape = size - target_hexagon = tvm.target.hexagon("v69") - func_tir = tvm.build( - sch.mod["main"], target=tvm.target.Target(target_hexagon, host=target_hexagon) - ) + func_tir = tvm.build(sch.mod["main"], target=get_hexagon_target("v69")) module = hexagon_session.load_module(func_tir) rng = default_rng() diff --git a/tests/python/contrib/test_hexagon/topi/test_add_subtract_multiply.py b/tests/python/contrib/test_hexagon/topi/test_add_subtract_multiply.py index fe70745143a9..711d725e842f 100755 --- a/tests/python/contrib/test_hexagon/topi/test_add_subtract_multiply.py +++ b/tests/python/contrib/test_hexagon/topi/test_add_subtract_multiply.py @@ -23,7 +23,12 @@ from tvm import te import tvm.topi.hexagon.slice_ops as sl import tvm.topi.hexagon.qnn as qn -from ..infrastructure import allocate_hexagon_array, transform_numpy, quantize_np +from ..infrastructure import ( + allocate_hexagon_array, + transform_numpy, + quantize_np, + get_hexagon_target, +) @tvm.testing.fixture @@ -285,7 +290,6 @@ def test_transform( op_name, ): output_shape = expected_output_np.shape - target_hexagon = tvm.target.hexagon("v69") A = te.placeholder(input_shape_A, name="A", dtype=dtype) B = te.placeholder(input_shape_B, name="B", dtype=dtype) if dtype == "float16": @@ -336,7 +340,7 @@ def test_transform( func = tvm.build( sch, [A, B, M], - tvm.target.Target(target_hexagon, host=target_hexagon), + get_hexagon_target("v69"), name="slice_op_with_transform", ) diff --git a/tests/python/contrib/test_hexagon/topi/test_argmax_slice.py b/tests/python/contrib/test_hexagon/topi/test_argmax_slice.py index 32d7a5097384..5ed86a1fcc92 100644 --- a/tests/python/contrib/test_hexagon/topi/test_argmax_slice.py +++ b/tests/python/contrib/test_hexagon/topi/test_argmax_slice.py @@ -22,7 +22,7 @@ from tvm import te import tvm.topi.hexagon.slice_ops as sl import tvm.contrib.hexagon -from ..infrastructure import allocate_hexagon_array, transform_numpy +from ..infrastructure import allocate_hexagon_array, transform_numpy, get_hexagon_target class TestArgMaxSlice: @@ -79,8 +79,6 @@ def test_argmax_slice( working_scope, ): """Top level testing function for argmax""" - target_hexagon = tvm.target.hexagon("v69") - target = tvm.target.Target(target_hexagon, host=target_hexagon) argmax_input = te.placeholder(input_shape, name="A", dtype=dtype) output = sl.argmax.argmax_compute(argmax_input, in_axis) argmax_func = te.create_prim_func([argmax_input, output]) @@ -101,7 +99,7 @@ def test_argmax_slice( with tvm.transform.PassContext(opt_level=3): tir_irm = tvm.lower(tir_s.mod, [argmax_input, output], name="argmax") runtime_module = tvm.build( - tir_irm, [argmax_input, output], target=target, name="argmax" + tir_irm, [argmax_input, output], target=get_hexagon_target("v69"), name="argmax" ) mod = hexagon_session.load_module(runtime_module) diff --git a/tests/python/contrib/test_hexagon/topi/test_avg_pool2d_slice.py b/tests/python/contrib/test_hexagon/topi/test_avg_pool2d_slice.py index 743519901542..6f6a7d762747 100644 --- a/tests/python/contrib/test_hexagon/topi/test_avg_pool2d_slice.py +++ b/tests/python/contrib/test_hexagon/topi/test_avg_pool2d_slice.py @@ -21,12 +21,15 @@ from tvm import te import tvm.testing -from tvm.topi import testing -from tvm.contrib.hexagon.build import HexagonLauncher from tvm.contrib.hexagon.session import Session import tvm.topi.hexagon.slice_ops as sl import tvm.topi.hexagon.qnn as qn -from ..infrastructure import allocate_hexagon_array, transform_numpy, quantize_np +from ..infrastructure import ( + allocate_hexagon_array, + transform_numpy, + quantize_np, + get_hexagon_target, +) from ..pytest_util import ( get_multitest_ids, create_populated_numpy_ndarray, @@ -401,13 +404,12 @@ def test_avg_pool2d_slice( schedule_args, hexagon_session: Session, ): - target_hexagon = tvm.target.hexagon("v69") in_data = transformed_input_np_padded with tvm.transform.PassContext(opt_level=3): func = tvm.build( *schedule_args, - tvm.target.Target(target_hexagon, host=target_hexagon), + get_hexagon_target("v69"), name="avg_pool2d", ) diff --git a/tests/python/contrib/test_hexagon/topi/test_batch_matmul.py b/tests/python/contrib/test_hexagon/topi/test_batch_matmul.py index c64477343943..d572bf1bb42b 100644 --- a/tests/python/contrib/test_hexagon/topi/test_batch_matmul.py +++ b/tests/python/contrib/test_hexagon/topi/test_batch_matmul.py @@ -27,6 +27,7 @@ from tvm.topi.utils import get_const_tuple from tvm.contrib.hexagon.session import Session +from ..infrastructure import get_hexagon_target dtype = tvm.testing.parameter( "float32", @@ -74,7 +75,7 @@ def get_ref_data(): func = tvm.build( s, [x, y, out], - tvm.target.Target(target_hexagon, host=target_hexagon), + get_hexagon_target("v68"), name="batch_matmul", ) mod = hexagon_session.load_module(func) @@ -125,7 +126,7 @@ def get_ref_data(): func = tvm.build( s, [x, y, out], - tvm.target.Target(target_hexagon, host=target_hexagon), + get_hexagon_target("v68"), name="batch_matmul_int8", ) mod = hexagon_session.load_module(func) diff --git a/tests/python/contrib/test_hexagon/topi/test_cast_slice.py b/tests/python/contrib/test_hexagon/topi/test_cast_slice.py index 1b235a4daf52..326370eb72d7 100644 --- a/tests/python/contrib/test_hexagon/topi/test_cast_slice.py +++ b/tests/python/contrib/test_hexagon/topi/test_cast_slice.py @@ -22,7 +22,7 @@ import tvm.testing from tvm import te import tvm.topi.hexagon.slice_ops as sl -from ..infrastructure import allocate_hexagon_array, transform_numpy +from ..infrastructure import allocate_hexagon_array, transform_numpy, get_hexagon_target class TestCastF16F32Slice2d: @@ -77,8 +77,6 @@ def test_cast_fp16_fp32_slice( if hexagon_session._launcher._serial_number != "simulator": pytest.skip(msg="Due to https://github.com/apache/tvm/issues/11957") - target_hexagon = tvm.target.hexagon("v69") - target = tvm.target.Target(target_hexagon, host=target_hexagon) cast_input = te.placeholder(input_shape, name="A", dtype=dtype) cast_output = sl.cast_f16_f32_compute(cast_input) cast_func = te.create_prim_func([cast_input, cast_output]) @@ -98,7 +96,9 @@ def test_cast_fp16_fp32_slice( ) with tvm.transform.PassContext(opt_level=3): tir_irm = tvm.lower(tir_s.mod, [cast_input, cast_output], name="cast_f16_f32") - runtime_module = tvm.build(tir_irm, target=target, name="cast_f16_f32") + runtime_module = tvm.build( + tir_irm, target=get_hexagon_target("v69"), name="cast_f16_f32" + ) mod = hexagon_session.load_module(runtime_module) mod(input_data, output_data) @@ -163,8 +163,6 @@ def test_cast_fp32_fp16_slice( if hexagon_session._launcher._serial_number != "simulator": pytest.skip(msg="Due to https://github.com/apache/tvm/issues/11957") - target_hexagon = tvm.target.hexagon("v69") - target = tvm.target.Target(target_hexagon, host=target_hexagon) cast_input = te.placeholder(input_shape, name="A", dtype=dtype) cast_output = sl.cast_f32_f16_compute(cast_input) cast_func = te.create_prim_func([cast_input, cast_output]) @@ -184,7 +182,9 @@ def test_cast_fp32_fp16_slice( ) with tvm.transform.PassContext(opt_level=3): tir_irm = tvm.lower(tir_s.mod, [cast_input, cast_output], name="cast_f32_f16") - runtime_module = tvm.build(tir_irm, target=target, name="cast_f32_f16") + runtime_module = tvm.build( + tir_irm, target=get_hexagon_target("v69"), name="cast_f32_f16" + ) mod = hexagon_session.load_module(runtime_module) mod(input_data, output_data) diff --git a/tests/python/contrib/test_hexagon/topi/test_clip.py b/tests/python/contrib/test_hexagon/topi/test_clip.py index ac6890171dba..3f8f5077c758 100755 --- a/tests/python/contrib/test_hexagon/topi/test_clip.py +++ b/tests/python/contrib/test_hexagon/topi/test_clip.py @@ -19,13 +19,11 @@ import numpy as np -from tvm import te, topi +from tvm import te import tvm.testing -from tvm.topi import testing -from tvm.contrib.hexagon.build import HexagonLauncher import tvm.topi.hexagon.slice_ops as sl -from ..infrastructure import allocate_hexagon_array, transform_numpy +from ..infrastructure import allocate_hexagon_array, transform_numpy, get_hexagon_target input_layout = tvm.testing.parameter( "nhwc-8h2w32c2w-2d", @@ -73,7 +71,6 @@ def test_clip_slice( hexagon_session, ): # establish target and input placeholder - target_hexagon = tvm.target.hexagon("v69") A = te.placeholder(input_shape, name="A", dtype=dtype) # get the compute function and schedule @@ -86,7 +83,7 @@ def test_clip_slice( with tvm.transform.PassContext(opt_level=3): func = tvm.build( tir_schedule.mod, - target=tvm.target.Target(target_hexagon, host=target_hexagon), + target=get_hexagon_target("v69"), name="clip", ) diff --git a/tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py b/tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py index e8efdb369590..5066a532df9b 100644 --- a/tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py +++ b/tests/python/contrib/test_hexagon/topi/test_conv2d_fp16_intrin.py @@ -23,6 +23,8 @@ import tvm.contrib.hexagon from tvm.topi.testing import conv2d_nhwc_python +from ..infrastructure import get_hexagon_target + def build_conv2d(target): """Build and the return the conv2d module that calls the intrinsic implementation""" @@ -199,9 +201,6 @@ def test_conv2d(self, act_shape, wgt_shape, inp_stride, inp_offset, hexagon_sess """Test conv2d intrinsic implementation""" assert act_shape[3] == wgt_shape[2] - target_hexagon = tvm.target.hexagon("v69") - target = tvm.target.Target(target_hexagon, host=target_hexagon) - # Currently, input offset does not affect the output shape def get_out_shape(ash, wsh, inp_stride): assert ash[3] == wsh[2] @@ -217,7 +216,7 @@ def get_out_shape(ash, wsh, inp_stride): act = np.random.rand(*act_shape).astype("float16") wgt = np.random.rand(*wgt_shape).astype("float16") - module = build_conv2d(target) + module = build_conv2d(get_hexagon_target("v68")) mod = hexagon_session.load_module(module) output = tvm.nd.array( diff --git a/tests/python/contrib/test_hexagon/topi/test_conv2d_nchw.py b/tests/python/contrib/test_hexagon/topi/test_conv2d_nchw.py index 01c20601b685..295b5458ee81 100644 --- a/tests/python/contrib/test_hexagon/topi/test_conv2d_nchw.py +++ b/tests/python/contrib/test_hexagon/topi/test_conv2d_nchw.py @@ -26,6 +26,7 @@ from tvm.topi.utils import get_const_tuple from tvm.topi.nn.utils import get_pad_tuple +from ..infrastructure import get_hexagon_target dtype = tvm.testing.parameter("float32") random_seed = tvm.testing.parameter(0) @@ -153,7 +154,7 @@ def test_conv2d_nchw( func = tvm.build( s, [A, W, bias, C], - tvm.target.Target(target_hexagon, host=target_hexagon), + get_hexagon_target("v68"), name=func_name, ) mod = hexagon_session.load_module(func) diff --git a/tests/python/contrib/test_hexagon/topi/test_conv2d_nhwc.py b/tests/python/contrib/test_hexagon/topi/test_conv2d_nhwc.py index 9acffff358e8..2d849a54a661 100644 --- a/tests/python/contrib/test_hexagon/topi/test_conv2d_nhwc.py +++ b/tests/python/contrib/test_hexagon/topi/test_conv2d_nhwc.py @@ -25,6 +25,8 @@ import tvm.topi.testing from tvm.topi.utils import get_const_tuple +from ..infrastructure import get_hexagon_target + dtype = tvm.testing.parameter("float32") @@ -81,9 +83,7 @@ def test_conv2d_nhwc( padding, dilation, ) - func = tvm.build( - s, [A, W, B], tvm.target.Target(target_hexagon, host=target_hexagon), name=func_name - ) + func = tvm.build(s, [A, W, B], get_hexagon_target("v68"), name=func_name) mod = hexagon_session.load_module(func) dev = hexagon_session.device diff --git a/tests/python/contrib/test_hexagon/topi/test_conv2d_slice.py b/tests/python/contrib/test_hexagon/topi/test_conv2d_slice.py index b3ef6890e290..242265169fb8 100755 --- a/tests/python/contrib/test_hexagon/topi/test_conv2d_slice.py +++ b/tests/python/contrib/test_hexagon/topi/test_conv2d_slice.py @@ -25,7 +25,7 @@ from tvm.topi.hexagon.slice_ops.conv2d import conv2d_compute, conv2d_schedule from tvm.topi.testing import conv2d_nhwc_python -from ..infrastructure import allocate_hexagon_array, transform_numpy +from ..infrastructure import allocate_hexagon_array, transform_numpy, get_hexagon_target input_layout = tvm.testing.parameter( "nhwc-8h2w32c2w-2d", @@ -286,9 +286,6 @@ def test_conv2d( input_tensor, weights, out_shape, stride, dilation, dtype, output_name ) - target_hexagon = tvm.target.hexagon("v69") - target = tvm.target.Target(target_hexagon, host=target_hexagon) - tir_schedule = conv2d_schedule( output_tensor, [input_tensor, weights], @@ -302,7 +299,7 @@ def test_conv2d( with tvm.transform.PassContext(opt_level=3): runtime_module = tvm.build( tir_schedule.mod, - target=target, + target=get_hexagon_target("v69"), name=func_name, ) diff --git a/tests/python/contrib/test_hexagon/topi/test_conv2d_transpose.py b/tests/python/contrib/test_hexagon/topi/test_conv2d_transpose.py index 8536603a3c20..1a5dc4a17824 100644 --- a/tests/python/contrib/test_hexagon/topi/test_conv2d_transpose.py +++ b/tests/python/contrib/test_hexagon/topi/test_conv2d_transpose.py @@ -16,6 +16,7 @@ # under the License. """Test code for transposed convolution.""" import numpy as np + import tvm from tvm.contrib.hexagon.session import Session import tvm.testing @@ -24,6 +25,7 @@ import tvm.topi.testing from tvm.topi.utils import get_const_tuple +from ..infrastructure import get_hexagon_target # TODO Should add kernal to tvm.testing.fixture @@ -131,8 +133,8 @@ def get_ref_data(): b = tvm.nd.array(np.zeros(get_const_tuple(B.shape), dtype=B.dtype), dev) c = tvm.nd.array(np.zeros(get_const_tuple(C.shape), dtype=C.dtype), dev) - func1 = tvm.build(s1, [A, W, B], tvm.target.Target(target_hexagon, host=target_hexagon)) - func2 = tvm.build(s2, [A, W, C], tvm.target.Target(target_hexagon, host=target_hexagon)) + func1 = tvm.build(s1, [A, W, B], get_hexagon_target("v68")) + func2 = tvm.build(s2, [A, W, C], get_hexagon_target("v68")) mod1 = hexagon_session.load_module(func1) mod2 = hexagon_session.load_module(func2) diff --git a/tests/python/contrib/test_hexagon/topi/test_dense.py b/tests/python/contrib/test_hexagon/topi/test_dense.py index 929108bb1492..b8f5741e803e 100644 --- a/tests/python/contrib/test_hexagon/topi/test_dense.py +++ b/tests/python/contrib/test_hexagon/topi/test_dense.py @@ -26,6 +26,8 @@ import tvm.topi.testing from tvm.topi.utils import get_const_tuple +from ..infrastructure import get_hexagon_target + random_seed = tvm.testing.parameter(0) use_bias = tvm.testing.parameter(True, False) @@ -100,9 +102,7 @@ def test_dense( D = topi.nn.relu(D) s = fschedule([D]) - func = tvm.build( - s, [A, B, C, D], tvm.target.Target(target_hexagon, host=target_hexagon), name="dense" - ) + func = tvm.build(s, [A, B, C, D], get_hexagon_target("v68"), name="dense") mod = hexagon_session.load_module(func) dev = hexagon_session.device diff --git a/tests/python/contrib/test_hexagon/topi/test_depth_to_space.py b/tests/python/contrib/test_hexagon/topi/test_depth_to_space.py index f74d13f641d5..3de9ec13497a 100644 --- a/tests/python/contrib/test_hexagon/topi/test_depth_to_space.py +++ b/tests/python/contrib/test_hexagon/topi/test_depth_to_space.py @@ -27,7 +27,7 @@ from tvm.topi.hexagon.slice_ops.depth_to_space import d2s_compute, d2s_schedule from tvm.topi.testing import depth_to_space_python -from ..infrastructure import allocate_hexagon_array, transform_numpy +from ..infrastructure import allocate_hexagon_array, transform_numpy, get_hexagon_target d2s_fp16_tests = ( @@ -97,9 +97,6 @@ def test_d2s_slice( Output = d2s_compute(Input, block_size, "NHWC", mode) - target_hexagon = tvm.target.hexagon("v69") - target = tvm.target.Target(target_hexagon, host=target_hexagon) - tir_s = d2s_schedule(Input, Output, input_layout, output_layout) input_data = allocate_hexagon_array( @@ -117,7 +114,7 @@ def test_d2s_slice( ) with tvm.transform.PassContext(opt_level=3): runtime_module = tvm.build( - tir_s.mod, [Input, Output], target=target, name="depth_to_space" + tir_s.mod, [Input, Output], target=get_hexagon_target("v69"), name="depth_to_space" ) mod = hexagon_session.load_module(runtime_module) diff --git a/tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py b/tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py index 5e09e691f743..048db97510d0 100644 --- a/tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py +++ b/tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py @@ -28,6 +28,7 @@ from tvm.topi.utils import get_const_tuple from tvm.topi.nn.utils import get_pad_tuple +from ..infrastructure import get_hexagon_target random_seed = tvm.testing.parameter(0) @@ -236,7 +237,7 @@ def test_conv2d( f = tvm.build( s, [Input, Filter, Scale, Shift, C], - tvm.target.Target(target_hexagon, host=target_hexagon), + get_hexagon_target("v68"), ) mod = hexagon_session.load_module(f) diff --git a/tests/python/contrib/test_hexagon/topi/test_dequantize_slice.py b/tests/python/contrib/test_hexagon/topi/test_dequantize_slice.py index e9b3dd132692..6ed217180aba 100644 --- a/tests/python/contrib/test_hexagon/topi/test_dequantize_slice.py +++ b/tests/python/contrib/test_hexagon/topi/test_dequantize_slice.py @@ -23,7 +23,12 @@ import tvm.testing from tvm import te from tvm.topi.hexagon import qnn -from ..infrastructure import allocate_hexagon_array, transform_numpy, quantize_np +from ..infrastructure import ( + allocate_hexagon_array, + transform_numpy, + quantize_np, + get_hexagon_target, +) class TestDequantizeSlice2d: @@ -78,8 +83,6 @@ def test_dequant_qnn( """ Top level testing function for dequantize """ - target_hexagon = tvm.target.hexagon("v69") - target = tvm.target.Target(target_hexagon, host=target_hexagon) dequant_input = te.placeholder(input_shape, name="A", dtype=dtype) @@ -104,7 +107,7 @@ def test_dequant_qnn( ) with tvm.transform.PassContext(opt_level=3): tir_irm = tvm.lower(tir_s.mod, [dequant_input, dequant_output], name="dequantize") - runtime_module = tvm.build(tir_irm, target=target, name="dequantize") + runtime_module = tvm.build(tir_irm, target=get_hexagon_target("v69"), name="dequantize") mod = hexagon_session.load_module(runtime_module) mod(input_data, output_data) diff --git a/tests/python/contrib/test_hexagon/topi/test_dwconv2d_slice.py b/tests/python/contrib/test_hexagon/topi/test_dwconv2d_slice.py index bc72e2bf4bbc..3e43718afd8d 100644 --- a/tests/python/contrib/test_hexagon/topi/test_dwconv2d_slice.py +++ b/tests/python/contrib/test_hexagon/topi/test_dwconv2d_slice.py @@ -26,7 +26,7 @@ from tvm.topi.testing import depthwise_conv2d_python_nhwc from tvm.topi.hexagon.slice_ops.dwconv2d import dwconv2d_compute, dwconv2d_schedule -from ..infrastructure import allocate_hexagon_array, transform_numpy +from ..infrastructure import allocate_hexagon_array, transform_numpy, get_hexagon_target @tvm.testing.fixture @@ -270,9 +270,6 @@ def test_dwconv2d( def transform_weights(height, width, in_channel, out_channel): return [out_channel // 32, height, width, in_channel, out_channel % 32] - target_hexagon = tvm.target.hexagon("v69") - target = tvm.target.Target(target_hexagon, host=target_hexagon) - tir_schedule = dwconv2d_schedule( output_tensor, [input_tensor, weights], in_out_layout, transform_weights ) @@ -281,7 +278,7 @@ def transform_weights(height, width, in_channel, out_channel): with tvm.transform.PassContext(opt_level=3, config={"tir.disable_assert": True}): runtime_module = tvm.build( tir_schedule.mod, - target=target, + target=get_hexagon_target("v69"), name=func_name, ) diff --git a/tests/python/contrib/test_hexagon/topi/test_max_pool2d_slice.py b/tests/python/contrib/test_hexagon/topi/test_max_pool2d_slice.py index de60ffc6df4d..f2ee76863cb6 100644 --- a/tests/python/contrib/test_hexagon/topi/test_max_pool2d_slice.py +++ b/tests/python/contrib/test_hexagon/topi/test_max_pool2d_slice.py @@ -21,14 +21,10 @@ from tvm import te import tvm.testing -from tvm.topi import testing -from tvm.contrib.hexagon.build import HexagonLauncher from tvm.contrib.hexagon.session import Session import tvm.topi.hexagon.slice_ops as sl -from ..infrastructure import allocate_hexagon_array, transform_numpy +from ..infrastructure import allocate_hexagon_array, transform_numpy, get_hexagon_target from ..pytest_util import ( - get_numpy_dtype_info, - get_test_id, get_multitest_ids, create_populated_numpy_ndarray, TensorContentRandom, @@ -341,7 +337,6 @@ def test_max_pool2d_slice( expected_output_np, hexagon_session: Session, ): - target_hexagon = tvm.target.hexagon("v69") A = te.placeholder(input_shape_padded, name="A", dtype=dtype) M = sl.max_pool2d_compute(A, output_shape, kernel, stride, dilation) @@ -365,7 +360,7 @@ def test_max_pool2d_slice( func = tvm.build( sch, [A, M], - tvm.target.Target(target_hexagon, host=target_hexagon), + get_hexagon_target("v69"), name="max_pool2d", ) diff --git a/tests/python/contrib/test_hexagon/topi/test_pad.py b/tests/python/contrib/test_hexagon/topi/test_pad.py index 631cb979dcbd..8a91eba471a2 100644 --- a/tests/python/contrib/test_hexagon/topi/test_pad.py +++ b/tests/python/contrib/test_hexagon/topi/test_pad.py @@ -22,6 +22,8 @@ from tvm.contrib.hexagon.session import Session from tvm.topi.utils import get_const_tuple +from ..infrastructure import get_hexagon_target + @tvm.testing.requires_hexagon def test_nn_pad(hexagon_session: Session): @@ -39,7 +41,7 @@ def test_nn_pad(hexagon_session: Session): fschedule = topi.hexagon.schedule_pad s = fschedule(C) - func = tvm.build(s, [A, C], tvm.target.Target(target_hexagon, host=target_hexagon), name="pad") + func = tvm.build(s, [A, C], get_hexagon_target("v68"), name="pad") mod = hexagon_session.load_module(func) dev = hexagon_session.device diff --git a/tests/python/contrib/test_hexagon/topi/test_pooling.py b/tests/python/contrib/test_hexagon/topi/test_pooling.py index 45e558e1b6dd..3bbf8657a012 100644 --- a/tests/python/contrib/test_hexagon/topi/test_pooling.py +++ b/tests/python/contrib/test_hexagon/topi/test_pooling.py @@ -25,6 +25,8 @@ import tvm.topi.testing from tvm.topi.utils import get_const_tuple +from ..infrastructure import get_hexagon_target + class TestAdaptivePool: dshape, out_size, pool_type, layout = tvm.testing.parameters( @@ -76,7 +78,7 @@ def test_adaptive_pool(self, hexagon_session: Session, dshape, out_size, pool_ty func = tvm.build( s, [data, out], - tvm.target.Target(target_hexagon, host=target_hexagon), + get_hexagon_target("v68"), name="adaptive-pool", ) mod = hexagon_session.load_module(func) @@ -171,7 +173,7 @@ def verify_poolnd( fschedule = topi.hexagon.schedule_pool s = fschedule(B, layout) - func = tvm.build(s, [A, B], tvm.target.Target(target_hexagon, host=target_hexagon), name="pool") + func = tvm.build(s, [A, B], get_hexagon_target("v68"), name="pool") mod = hexagon_session.load_module(func) dev = hexagon_session.device diff --git a/tests/python/contrib/test_hexagon/topi/test_quantize.py b/tests/python/contrib/test_hexagon/topi/test_quantize.py index 2c1718d29465..0b6e1dfa0e73 100755 --- a/tests/python/contrib/test_hexagon/topi/test_quantize.py +++ b/tests/python/contrib/test_hexagon/topi/test_quantize.py @@ -14,13 +14,17 @@ # KIND, either express or implied. See the License for the # specific language governing permissions and limitations # under the License. -import pytest import numpy as np import tvm from tvm import te import tvm.topi.hexagon.qnn as s1 -from ..infrastructure import allocate_hexagon_array, transform_numpy, quantize_np +from ..infrastructure import ( + allocate_hexagon_array, + transform_numpy, + quantize_np, + get_hexagon_target, +) @tvm.testing.fixture @@ -70,7 +74,6 @@ def test_quantize( output_layout, hexagon_session, ): - target_hexagon = tvm.target.hexagon("v69") A = te.placeholder(input_shape, name="A", dtype=input_dtype) M = s1.quantize_compute(A, scale, zero_point, output_dtype) @@ -86,7 +89,7 @@ def test_quantize( func = tvm.build( sch, [A, M], - tvm.target.Target(target_hexagon, host=target_hexagon), + get_hexagon_target("v69"), name="quantize", ) diff --git a/tests/python/contrib/test_hexagon/topi/test_reduce.py b/tests/python/contrib/test_hexagon/topi/test_reduce.py index a844e1d51206..668431e7ba49 100644 --- a/tests/python/contrib/test_hexagon/topi/test_reduce.py +++ b/tests/python/contrib/test_hexagon/topi/test_reduce.py @@ -22,6 +22,7 @@ from tvm import te from tvm.contrib.hexagon.session import Session +from ..infrastructure import get_hexagon_target in_shape, axis, keepdims, reduce_type, dtype = tvm.testing.parameters( ((32,), 0, False, "argmax", "float32"), @@ -130,9 +131,7 @@ def test_reduce_map( fschedule = topi.hexagon.schedule_reduce s = fschedule(B) - func = tvm.build( - s, [A, B], tvm.target.Target(target_hexagon, host=target_hexagon), name=reduce_type - ) + func = tvm.build(s, [A, B], get_hexagon_target("v68"), name=reduce_type) mod = hexagon_session.load_module(func) dev = hexagon_session.device diff --git a/tests/python/contrib/test_hexagon/topi/test_relu_slice.py b/tests/python/contrib/test_hexagon/topi/test_relu_slice.py index 87978f586541..fd04cca061da 100644 --- a/tests/python/contrib/test_hexagon/topi/test_relu_slice.py +++ b/tests/python/contrib/test_hexagon/topi/test_relu_slice.py @@ -22,9 +22,8 @@ import tvm.testing from tvm.topi.hexagon.slice_ops.relu import relu_compute, relu_stir_schedule from tvm import te -from tvm.contrib.hexagon.build import HexagonLauncher -from ..infrastructure import allocate_hexagon_array, transform_numpy +from ..infrastructure import allocate_hexagon_array, transform_numpy, get_hexagon_target @tvm.testing.fixture @@ -82,9 +81,6 @@ def test_relu( OutputTensor = relu_compute(InputTensor) - target_hexagon = tvm.target.hexagon("v69") - target = tvm.target.Target(target_hexagon, host=target_hexagon) - tir_s = relu_stir_schedule(InputTensor, OutputTensor, input_layout, output_layout) input_data = allocate_hexagon_array( @@ -103,7 +99,7 @@ def test_relu( func_name = "relu" with tvm.transform.PassContext(opt_level=3): - runtime_module = tvm.build(tir_s.mod, target=target, name=func_name) + runtime_module = tvm.build(tir_s.mod, target=get_hexagon_target("v69"), name=func_name) mod = hexagon_session.load_module(runtime_module) diff --git a/tests/python/contrib/test_hexagon/topi/test_reshape.py b/tests/python/contrib/test_hexagon/topi/test_reshape.py index 7df29a02abff..38b8a9cf9a82 100644 --- a/tests/python/contrib/test_hexagon/topi/test_reshape.py +++ b/tests/python/contrib/test_hexagon/topi/test_reshape.py @@ -21,11 +21,9 @@ import tvm import tvm.testing import tvm.topi.hexagon.slice_ops as sl -from tvm import te, topi -from tvm.contrib.hexagon.build import HexagonLauncher -from tvm.topi import testing +from tvm import te -from ..infrastructure import allocate_hexagon_array, transform_numpy +from ..infrastructure import allocate_hexagon_array, transform_numpy, get_hexagon_target def reshape_helper( @@ -40,8 +38,6 @@ def reshape_helper( hexagon_session, ): - target_hexagon = tvm.target.hexagon("v69") - target = tvm.target.Target(target_hexagon, host=target_hexagon) A = te.placeholder(input_shape, name="A", dtype=data_type) if func == "reshape": D = fcompute(A, output_shape) @@ -56,7 +52,7 @@ def reshape_helper( input_layout, ) with tvm.transform.PassContext(opt_level=3): - runtime_module = tvm.build(tir_s.mod, target=target, name=func) + runtime_module = tvm.build(tir_s.mod, target=get_hexagon_target("v69"), name=func) mod = hexagon_session.load_module(runtime_module) diff --git a/tests/python/contrib/test_hexagon/topi/test_resize2d.py b/tests/python/contrib/test_hexagon/topi/test_resize2d.py index 1ef9f50977c5..80cfba5c6c9e 100755 --- a/tests/python/contrib/test_hexagon/topi/test_resize2d.py +++ b/tests/python/contrib/test_hexagon/topi/test_resize2d.py @@ -21,7 +21,7 @@ from tvm import te from tvm.topi.testing import resize2d_python import tvm.topi.hexagon as s1 -from ..infrastructure import allocate_hexagon_array, transform_numpy +from ..infrastructure import allocate_hexagon_array, transform_numpy, get_hexagon_target @tvm.testing.fixture @@ -123,7 +123,6 @@ def test_resize2d( method, hexagon_session, ): - target_hexagon = tvm.target.hexagon("v69") A = te.placeholder(input_shape, name="A", dtype=dtype) M = s1.resize2d_compute( @@ -153,7 +152,7 @@ def test_resize2d( func = tvm.build( sch, [A, M], - tvm.target.Target(target_hexagon, host=target_hexagon), + get_hexagon_target("v69"), name="resize2d", ) diff --git a/tests/python/contrib/test_hexagon/topi/test_softmax.py b/tests/python/contrib/test_hexagon/topi/test_softmax.py index d1c78842b5ff..4fd981407836 100644 --- a/tests/python/contrib/test_hexagon/topi/test_softmax.py +++ b/tests/python/contrib/test_hexagon/topi/test_softmax.py @@ -26,6 +26,8 @@ import tvm.topi.testing from tvm.topi.utils import get_const_tuple +from ..infrastructure import get_hexagon_target + dtype = tvm.testing.parameter( "float16", "float32", @@ -83,9 +85,7 @@ def get_ref_data(shape): fschedule = topi.hexagon.schedule_softmax s = fschedule(B) - func = tvm.build( - s, [A, B], tvm.target.Target(target_hexagon, host=target_hexagon), name="softmax" - ) + func = tvm.build(s, [A, B], get_hexagon_target("v68"), name="softmax") mod = hexagon_session.load_module(func) dev = hexagon_session.device diff --git a/tests/python/contrib/test_hexagon/topi/test_softmax_slice.py b/tests/python/contrib/test_hexagon/topi/test_softmax_slice.py index 91b51cb5cc75..1329fda7aa4a 100644 --- a/tests/python/contrib/test_hexagon/topi/test_softmax_slice.py +++ b/tests/python/contrib/test_hexagon/topi/test_softmax_slice.py @@ -20,7 +20,8 @@ from tvm import te from tvm.topi.testing import softmax_python import tvm.topi.hexagon.slice_ops as sl -from ..infrastructure import allocate_hexagon_array + +from ..infrastructure import allocate_hexagon_array, get_hexagon_target def transform_numpy(arr_np, layout): diff --git a/tests/python/contrib/test_hexagon/topi/test_tanh_slice.py b/tests/python/contrib/test_hexagon/topi/test_tanh_slice.py index d488d7dd46dd..02c587b9809c 100644 --- a/tests/python/contrib/test_hexagon/topi/test_tanh_slice.py +++ b/tests/python/contrib/test_hexagon/topi/test_tanh_slice.py @@ -23,7 +23,7 @@ from tvm import te import tvm.topi.hexagon.slice_ops as sl import tvm.contrib.hexagon -from ..infrastructure import allocate_hexagon_array, transform_numpy +from ..infrastructure import allocate_hexagon_array, transform_numpy, get_hexagon_target # pylint: disable=invalid-name @@ -71,8 +71,6 @@ def test_tanh( ): """Top Level testing function for tanh fp16 op""" - target_hexagon = tvm.target.hexagon("v69") - target = tvm.target.Target(target_hexagon, host=target_hexagon) A = te.placeholder(input_shape, name="A", dtype=dtype) M = sl.tanh_te_compute(A) tanhf16_func = te.create_prim_func([A, M]) @@ -92,7 +90,7 @@ def test_tanh( ) with tvm.transform.PassContext(opt_level=3): tir_irm = tvm.lower(tir_s.mod, [A, M], name="tanhf16") - runtime_module = tvm.build(tir_irm, target=target, name="tanhf16") + runtime_module = tvm.build(tir_irm, target=get_hexagon_target("v69"), name="tanhf16") mod = hexagon_session.load_module(runtime_module) mod(A_data, M_data) From 15a69ad560c2be6a73a714d3d4fc4d72f2f7ae48 Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Tue, 4 Oct 2022 10:44:08 -0700 Subject: [PATCH 3/9] refactor target --- tests/python/contrib/test_hexagon/topi/test_batch_matmul.py | 6 ++---- tests/python/contrib/test_hexagon/topi/test_conv2d_nchw.py | 3 +-- tests/python/contrib/test_hexagon/topi/test_conv2d_nhwc.py | 4 +--- .../contrib/test_hexagon/topi/test_conv2d_transpose.py | 5 +---- tests/python/contrib/test_hexagon/topi/test_dense.py | 3 +-- .../contrib/test_hexagon/topi/test_depthwise_conv2d.py | 4 +--- tests/python/contrib/test_hexagon/topi/test_pad.py | 3 +-- tests/python/contrib/test_hexagon/topi/test_pooling.py | 6 ++---- tests/python/contrib/test_hexagon/topi/test_reduce.py | 3 +-- tests/python/contrib/test_hexagon/topi/test_softmax.py | 3 +-- 10 files changed, 12 insertions(+), 28 deletions(-) diff --git a/tests/python/contrib/test_hexagon/topi/test_batch_matmul.py b/tests/python/contrib/test_hexagon/topi/test_batch_matmul.py index d572bf1bb42b..f3273ea8b65b 100644 --- a/tests/python/contrib/test_hexagon/topi/test_batch_matmul.py +++ b/tests/python/contrib/test_hexagon/topi/test_batch_matmul.py @@ -64,8 +64,7 @@ def get_ref_data(): # get the test data a_np, b_np, c_np = get_ref_data() - target_hexagon = tvm.target.hexagon("v68") - with tvm.target.Target(target_hexagon): + with tvm.target.Target(get_hexagon_target("v68")): fcompute = topi.nn.batch_matmul fschedule = topi.hexagon.schedule_batch_matmul out = fcompute(x, y) @@ -116,8 +115,7 @@ def get_ref_data(): # get the test data a_np, b_np, c_np = get_ref_data() - target_hexagon = tvm.target.hexagon("v68") - with tvm.target.Target(target_hexagon): + with tvm.target.Target(get_hexagon_target("v68")): fcompute = topi.nn.batch_matmul fschedule = topi.hexagon.schedule_batch_matmul out = fcompute(x, y) diff --git a/tests/python/contrib/test_hexagon/topi/test_conv2d_nchw.py b/tests/python/contrib/test_hexagon/topi/test_conv2d_nchw.py index 295b5458ee81..0b94d6e781a7 100644 --- a/tests/python/contrib/test_hexagon/topi/test_conv2d_nchw.py +++ b/tests/python/contrib/test_hexagon/topi/test_conv2d_nchw.py @@ -106,7 +106,6 @@ def test_conv2d_nchw( add_bias, apply_relu, ): - target_hexagon = tvm.target.hexagon("v68") pad_top, pad_left, pad_bottom, pad_right = get_pad_tuple(padding, (kernel, kernel)) padding_sum = pad_top + pad_left + pad_bottom + pad_right @@ -130,7 +129,7 @@ def test_conv2d_nchw( gap_size = np.nextafter(c_np.max(), np.inf, dtype=c_np.dtype) - c_np.max() tol = {"rtol": 1e-3, "atol": num_values_summed * gap_size / 2} - with tvm.target.Target(target_hexagon): + with tvm.target.Target(get_hexagon_target("v68")): fcompute = topi.nn.conv2d_nchw fschedule = topi.hexagon.schedule_conv2d_nchw C = fcompute(A, W, (stride, stride), padding, (dilation, dilation), dtype) diff --git a/tests/python/contrib/test_hexagon/topi/test_conv2d_nhwc.py b/tests/python/contrib/test_hexagon/topi/test_conv2d_nhwc.py index 2d849a54a661..2068f1e6e6fc 100644 --- a/tests/python/contrib/test_hexagon/topi/test_conv2d_nhwc.py +++ b/tests/python/contrib/test_hexagon/topi/test_conv2d_nhwc.py @@ -59,14 +59,12 @@ def test_conv2d_nhwc( padding, dilation, ): - target_hexagon = tvm.target.hexagon("v68") - a_np, w_np, b_np = ref_data A = te.placeholder(a_np.shape, name="A", dtype=dtype) W = te.placeholder(w_np.shape, name="W", dtype=dtype) - with tvm.target.Target(target_hexagon): + with tvm.target.Target(get_hexagon_target("v68")): fcompute = topi.nn.conv2d_nhwc fschedule = topi.hexagon.schedule_conv2d_nhwc B = fcompute(A, W, stride, padding, dilation, dtype) diff --git a/tests/python/contrib/test_hexagon/topi/test_conv2d_transpose.py b/tests/python/contrib/test_hexagon/topi/test_conv2d_transpose.py index 1a5dc4a17824..40c8efa1cec2 100644 --- a/tests/python/contrib/test_hexagon/topi/test_conv2d_transpose.py +++ b/tests/python/contrib/test_hexagon/topi/test_conv2d_transpose.py @@ -81,9 +81,6 @@ def test_conv2d( output_padding, random_seed, ): - - target_hexagon = tvm.target.hexagon("v68") - in_height, in_width = in_size kernel_height, kernel_width = (1, 1) stride_height, stride_width = stride @@ -118,7 +115,7 @@ def get_ref_data(): output_padding, ) - with tvm.target.Target(target_hexagon): + with tvm.target.Target(get_hexagon_target("v68")): fcompute = topi.nn.conv2d_transpose_nchw fschedule = topi.hexagon.schedule_conv2d_transpose_nchw B = fcompute(*fcompute_args) diff --git a/tests/python/contrib/test_hexagon/topi/test_dense.py b/tests/python/contrib/test_hexagon/topi/test_dense.py index b8f5741e803e..c76006ac08c2 100644 --- a/tests/python/contrib/test_hexagon/topi/test_dense.py +++ b/tests/python/contrib/test_hexagon/topi/test_dense.py @@ -96,8 +96,7 @@ def test_dense( fcompute = topi.nn.dense fschedule = topi.hexagon.schedule_dense - target_hexagon = tvm.target.hexagon("v68") - with tvm.target.Target(target_hexagon): + with tvm.target.Target(get_hexagon_target("v68")): D = fcompute(A, B, C if use_bias else None, out_dtype) D = topi.nn.relu(D) s = fschedule([D]) diff --git a/tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py b/tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py index 048db97510d0..063541cc21a0 100644 --- a/tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py +++ b/tests/python/contrib/test_hexagon/topi/test_depthwise_conv2d.py @@ -176,8 +176,6 @@ def test_conv2d( dilation, ref_data, ): - target_hexagon = tvm.target.hexagon("v68") - # Transform the padding argument from 'str' to 'tuple' to # match the "workload" tuple in TopHub. Which padding_args to # use for each layout chosen to reproduce previous behavior. @@ -217,7 +215,7 @@ def test_conv2d( out_dtype, ) - with tvm.target.Target(target_hexagon): + with tvm.target.Target(get_hexagon_target("v68")): # Declare, build schedule if layout == "NCHW": fcompute = topi.nn.depthwise_conv2d_nchw diff --git a/tests/python/contrib/test_hexagon/topi/test_pad.py b/tests/python/contrib/test_hexagon/topi/test_pad.py index 8a91eba471a2..06b939bf6409 100644 --- a/tests/python/contrib/test_hexagon/topi/test_pad.py +++ b/tests/python/contrib/test_hexagon/topi/test_pad.py @@ -36,8 +36,7 @@ def test_nn_pad(hexagon_session: Session): C = topi.nn.pad(A, [0, 1, 1, 0], [0, 1, 1, 0], pad_value=0) - target_hexagon = tvm.target.hexagon("v68") - with tvm.target.Target(target_hexagon): + with tvm.target.Target(get_hexagon_target("v68")): fschedule = topi.hexagon.schedule_pad s = fschedule(C) diff --git a/tests/python/contrib/test_hexagon/topi/test_pooling.py b/tests/python/contrib/test_hexagon/topi/test_pooling.py index 3bbf8657a012..ecc998875296 100644 --- a/tests/python/contrib/test_hexagon/topi/test_pooling.py +++ b/tests/python/contrib/test_hexagon/topi/test_pooling.py @@ -70,8 +70,7 @@ def test_adaptive_pool(self, hexagon_session: Session, dshape, out_size, pool_ty assert len(out_size) == 3 out = topi.nn.adaptive_pool3d(data, out_size, pool_type, layout) - target_hexagon = tvm.target.hexagon("v68") - with tvm.target.Target(target_hexagon): + with tvm.target.Target(get_hexagon_target("v68")): fschedule = topi.hexagon.schedule_adaptive_pool s = fschedule(out) @@ -168,8 +167,7 @@ def verify_poolnd( np.testing.assert_equal(tuple(output_shape), tuple(ref_np.shape)) - target_hexagon = tvm.target.hexagon("v68") - with tvm.target.Target(target_hexagon): + with tvm.target.Target(get_hexagon_target("v68")): fschedule = topi.hexagon.schedule_pool s = fschedule(B, layout) diff --git a/tests/python/contrib/test_hexagon/topi/test_reduce.py b/tests/python/contrib/test_hexagon/topi/test_reduce.py index 668431e7ba49..8fc0b6d901ab 100644 --- a/tests/python/contrib/test_hexagon/topi/test_reduce.py +++ b/tests/python/contrib/test_hexagon/topi/test_reduce.py @@ -126,8 +126,7 @@ def test_reduce_map( else: raise NotImplementedError - target_hexagon = tvm.target.hexagon("v68") - with tvm.target.Target(target_hexagon): + with tvm.target.Target(get_hexagon_target("v68")): fschedule = topi.hexagon.schedule_reduce s = fschedule(B) diff --git a/tests/python/contrib/test_hexagon/topi/test_softmax.py b/tests/python/contrib/test_hexagon/topi/test_softmax.py index 4fd981407836..91f348494d6d 100644 --- a/tests/python/contrib/test_hexagon/topi/test_softmax.py +++ b/tests/python/contrib/test_hexagon/topi/test_softmax.py @@ -80,8 +80,7 @@ def get_ref_data(shape): # get the test data a_np, b_np = get_ref_data(shape) - target_hexagon = tvm.target.hexagon("v68") - with tvm.target.Target(target_hexagon): + with tvm.target.Target(get_hexagon_target("v68")): fschedule = topi.hexagon.schedule_softmax s = fschedule(B) From 9dcbf0d6c6492a099ceab50128fbf144379aa184 Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Tue, 4 Oct 2022 10:47:22 -0700 Subject: [PATCH 4/9] fix permission --- tests/python/contrib/test_hexagon/test_2d_physical_buffers.py | 0 1 file changed, 0 insertions(+), 0 deletions(-) mode change 100755 => 100644 tests/python/contrib/test_hexagon/test_2d_physical_buffers.py diff --git a/tests/python/contrib/test_hexagon/test_2d_physical_buffers.py b/tests/python/contrib/test_hexagon/test_2d_physical_buffers.py old mode 100755 new mode 100644 From 280011ae5ed385b85efa329040ab3f4301561f0e Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Tue, 4 Oct 2022 10:51:59 -0700 Subject: [PATCH 5/9] cleanup --- .../test_hexagon/test_benchmark_elemwise_add.py | 9 +++------ .../python/contrib/test_hexagon/test_memory_alloc.py | 11 +++-------- 2 files changed, 6 insertions(+), 14 deletions(-) diff --git a/tests/python/contrib/test_hexagon/test_benchmark_elemwise_add.py b/tests/python/contrib/test_hexagon/test_benchmark_elemwise_add.py index b15219ebc00e..7ccdfe3150e4 100644 --- a/tests/python/contrib/test_hexagon/test_benchmark_elemwise_add.py +++ b/tests/python/contrib/test_hexagon/test_benchmark_elemwise_add.py @@ -20,25 +20,22 @@ import os.path import sys import tempfile - import numpy as np import pytest + import tvm.script import tvm.testing from tvm.contrib.hexagon.build import HexagonLauncherRPC from tvm.script import tir as T from . import benchmark_util as bu +from .infrastructure import get_hexagon_target _SHOULD_SKIP_BENCHMARKS, _SKIP_BENCHMARKS_REASON = bu.skip_bencharks_flag_and_reason() # This is a fixed detail of the v68 architecture. HVX_VECTOR_BYTES = 128 -_HEXAGON_TARGET = tvm.target.hexagon("v69", link_params=True) - -_SUPER_TARGET = tvm.target.Target(_HEXAGON_TARGET, host=_HEXAGON_TARGET) - # NOTE on server ports: # These tests use different port numbers for the RPC server (7070 + ...). # The reason is that an RPC session cannot be gracefully closed without @@ -219,7 +216,7 @@ def _benchmark_hexagon_elementwise_add_kernel( input2, output, ], - _SUPER_TARGET, + get_hexagon_target("v69"), name=_PRIMFUNC_NAME, ) diff --git a/tests/python/contrib/test_hexagon/test_memory_alloc.py b/tests/python/contrib/test_hexagon/test_memory_alloc.py index fd948ea524f2..a6d011eddd5a 100644 --- a/tests/python/contrib/test_hexagon/test_memory_alloc.py +++ b/tests/python/contrib/test_hexagon/test_memory_alloc.py @@ -15,20 +15,14 @@ # specific language governing permissions and limitations # under the License. -import os import os.path -import sys -import tempfile import numpy as np -import pytest import tvm from tvm.script import tir as T -from .infrastructure import allocate_hexagon_array - -_HEXAGON_TARGET = tvm.target.hexagon("v69", link_params=True) +from .infrastructure import allocate_hexagon_array, get_hexagon_target @tvm.testing.fixture @@ -63,7 +57,8 @@ def test_global_axis_separator( self, hexagon_session, generated_func, shape, dtype, scope, axis_separators ): mod1 = tvm.build( - generated_func, target=tvm.target.Target(_HEXAGON_TARGET, host=_HEXAGON_TARGET) + generated_func, + target=get_hexagon_target("v69"), ) mod2 = hexagon_session.load_module(mod1) From cdff61402c52d4b981daa65dbb435085eed5e17f Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Tue, 4 Oct 2022 14:32:25 -0700 Subject: [PATCH 6/9] fix target --- tests/python/contrib/test_hexagon/test_parallel_hvx.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/python/contrib/test_hexagon/test_parallel_hvx.py b/tests/python/contrib/test_hexagon/test_parallel_hvx.py index 75ba885c8fff..6ebe03d4e6b1 100644 --- a/tests/python/contrib/test_hexagon/test_parallel_hvx.py +++ b/tests/python/contrib/test_hexagon/test_parallel_hvx.py @@ -134,7 +134,7 @@ def operator(a: T.handle, b: T.handle, c: T.handle) -> None: def evaluate(hexagon_session, shape_dtypes, expected_output_producer, sch): a_shape, a_dtype, b_shape, b_dtype, c_shape, c_dtype = shape_dtypes - func_tir = tvm.build(sch.mod["main"], get_hexagon_target("v68")) + func_tir = tvm.build(sch.mod["main"], target=get_hexagon_target("v68")) module = hexagon_session.load_module(func_tir) rng = default_rng() From 05bc6ee0d4b884fefa052b3732170e91006fbd2c Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Tue, 4 Oct 2022 15:10:57 -0700 Subject: [PATCH 7/9] remove target fixture from test_2d_physical_buffers --- .../test_hexagon/test_2d_physical_buffers.py | 16 +++------------- 1 file changed, 3 insertions(+), 13 deletions(-) diff --git a/tests/python/contrib/test_hexagon/test_2d_physical_buffers.py b/tests/python/contrib/test_hexagon/test_2d_physical_buffers.py index cba6ddc4433a..7804ae2e4898 100644 --- a/tests/python/contrib/test_hexagon/test_2d_physical_buffers.py +++ b/tests/python/contrib/test_hexagon/test_2d_physical_buffers.py @@ -32,7 +32,7 @@ from tvm.contrib.hexagon.pytest_plugin import requires_hexagon_toolchain from tvm.tir.stmt_functor import post_order_visit -from .infrastructure import allocate_hexagon_array +from .infrastructure import allocate_hexagon_array, get_hexagon_target # Disabling invalid name as pylint assumes global variables as constants and # expects them to be all upper-case. Since these are used as @@ -79,19 +79,9 @@ @tvm.testing.fixture -def target_host(target): +def target_host(): """Return tvm target.Target with host attached""" - target = tvm.target.Target(target) - - if target.kind.name == "hexagon": - # Shouldn't have to modify the target here, current - # workaround. In the future, should move the parameter - # handling from tvm.target to target_kind.cc. - target = tvm.target.hexagon("v68", link_params=True) - host = target - else: - host = None - return tvm.target.Target(target, host=host) + return get_hexagon_target("v68") # Disabling redefined-outer-name for the whole file as there isn't any easy From 510fbfb3d3157672e055e00e1a2a14bf5558435a Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Tue, 4 Oct 2022 16:15:29 -0700 Subject: [PATCH 8/9] fix target fixture in test_hexagon/conv2d tests --- .../python/contrib/test_hexagon/conv2d/test_conv2d_blocked.py | 3 +-- .../python/contrib/test_hexagon/conv2d/test_conv2d_conv2d.py | 4 ++-- tests/python/contrib/test_hexagon/infrastructure.py | 2 +- 3 files changed, 4 insertions(+), 5 deletions(-) diff --git a/tests/python/contrib/test_hexagon/conv2d/test_conv2d_blocked.py b/tests/python/contrib/test_hexagon/conv2d/test_conv2d_blocked.py index 07f6c2613dbc..31bcc027681b 100644 --- a/tests/python/contrib/test_hexagon/conv2d/test_conv2d_blocked.py +++ b/tests/python/contrib/test_hexagon/conv2d/test_conv2d_blocked.py @@ -150,7 +150,6 @@ class BaseConv2d: class TestConv2dPackedFilter(BaseConv2d): """Conv2d packed filter test class""" - @tvm.testing.parametrize_targets("llvm") @tvm.testing.skip_if_32bit(reason="Test known to be flaky on i386 machines") def test_conv2d( self, @@ -164,11 +163,11 @@ def test_conv2d( k_split_factor, h_split_factor, dtype, - target, ): """conv2d test""" # TODO: no support for dilation dilation = 1 + target = "llvm" shape_input = [batch, in_size, in_size, in_channel] shape_filter_oihw = [out_channel, in_channel, kernel_size, kernel_size] diff --git a/tests/python/contrib/test_hexagon/conv2d/test_conv2d_conv2d.py b/tests/python/contrib/test_hexagon/conv2d/test_conv2d_conv2d.py index fa770c9be313..319c1e06f1d8 100644 --- a/tests/python/contrib/test_hexagon/conv2d/test_conv2d_conv2d.py +++ b/tests/python/contrib/test_hexagon/conv2d/test_conv2d_conv2d.py @@ -173,7 +173,6 @@ class BaseConv2dConv2d: class TestConv2dConv2dPackedFilter(BaseConv2dConv2d): """Conv2d-Conv2d packed filter test class""" - @tvm.testing.parametrize_targets("llvm") @tvm.testing.skip_if_32bit(reason="Test known to be flaky on i386 machines") def test_conv2d( self, @@ -190,12 +189,13 @@ def test_conv2d( k_split_factor, h_split_factor, dtype, - target, ): """conv2d-conv2d test""" # TODO: no support for padding in conv2d #2 pad2 = 0 + target = "llvm" + # TODO: no support for dilation dilation1 = 1 dilation2 = 1 diff --git a/tests/python/contrib/test_hexagon/infrastructure.py b/tests/python/contrib/test_hexagon/infrastructure.py index 1d854b5fd64f..6f7e1904da2f 100644 --- a/tests/python/contrib/test_hexagon/infrastructure.py +++ b/tests/python/contrib/test_hexagon/infrastructure.py @@ -128,7 +128,7 @@ def get_packed_filter_shape(logical_shape_oihw): return physical_shape_oihw8i32o4i -def build_and_run(inputs, func, target, target_host, *args, **kwargs): +def build_and_run(inputs, func, target: str, target_host: str, *args, **kwargs): """build and run the function func""" schedule, placeholders, binds = func(*args, **kwargs) From 51065b0cecabd8014d82e830078fa76aec0c31ed Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Wed, 5 Oct 2022 10:30:57 -0700 Subject: [PATCH 9/9] address comments --- .../python/contrib/test_hexagon/conv2d/test_conv2d_blocked.py | 3 ++- .../python/contrib/test_hexagon/conv2d/test_conv2d_conv2d.py | 4 ++-- .../contrib/test_hexagon/test_benchmark_elemwise_add.py | 1 + 3 files changed, 5 insertions(+), 3 deletions(-) diff --git a/tests/python/contrib/test_hexagon/conv2d/test_conv2d_blocked.py b/tests/python/contrib/test_hexagon/conv2d/test_conv2d_blocked.py index 31bcc027681b..07f6c2613dbc 100644 --- a/tests/python/contrib/test_hexagon/conv2d/test_conv2d_blocked.py +++ b/tests/python/contrib/test_hexagon/conv2d/test_conv2d_blocked.py @@ -150,6 +150,7 @@ class BaseConv2d: class TestConv2dPackedFilter(BaseConv2d): """Conv2d packed filter test class""" + @tvm.testing.parametrize_targets("llvm") @tvm.testing.skip_if_32bit(reason="Test known to be flaky on i386 machines") def test_conv2d( self, @@ -163,11 +164,11 @@ def test_conv2d( k_split_factor, h_split_factor, dtype, + target, ): """conv2d test""" # TODO: no support for dilation dilation = 1 - target = "llvm" shape_input = [batch, in_size, in_size, in_channel] shape_filter_oihw = [out_channel, in_channel, kernel_size, kernel_size] diff --git a/tests/python/contrib/test_hexagon/conv2d/test_conv2d_conv2d.py b/tests/python/contrib/test_hexagon/conv2d/test_conv2d_conv2d.py index 319c1e06f1d8..fa770c9be313 100644 --- a/tests/python/contrib/test_hexagon/conv2d/test_conv2d_conv2d.py +++ b/tests/python/contrib/test_hexagon/conv2d/test_conv2d_conv2d.py @@ -173,6 +173,7 @@ class BaseConv2dConv2d: class TestConv2dConv2dPackedFilter(BaseConv2dConv2d): """Conv2d-Conv2d packed filter test class""" + @tvm.testing.parametrize_targets("llvm") @tvm.testing.skip_if_32bit(reason="Test known to be flaky on i386 machines") def test_conv2d( self, @@ -189,13 +190,12 @@ def test_conv2d( k_split_factor, h_split_factor, dtype, + target, ): """conv2d-conv2d test""" # TODO: no support for padding in conv2d #2 pad2 = 0 - target = "llvm" - # TODO: no support for dilation dilation1 = 1 dilation2 = 1 diff --git a/tests/python/contrib/test_hexagon/test_benchmark_elemwise_add.py b/tests/python/contrib/test_hexagon/test_benchmark_elemwise_add.py index 7ccdfe3150e4..3dcb9a880e00 100644 --- a/tests/python/contrib/test_hexagon/test_benchmark_elemwise_add.py +++ b/tests/python/contrib/test_hexagon/test_benchmark_elemwise_add.py @@ -20,6 +20,7 @@ import os.path import sys import tempfile + import numpy as np import pytest