[Hexagon] Do not use target test fixture in Hexagon tests

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)
 
@@ -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)

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

tests/python/contrib/test_hexagon/test_autotvm.py

@@ -19,15 +19,15 @@
 
 import contextlib
 import os
-import sys
-
 import pytest
 
 import tvm
 import tvm.testing
 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()

tests/python/contrib/test_hexagon/test_benchmark_elemwise_add.py

@@ -23,22 +23,20 @@
 
 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 +217,7 @@ def _benchmark_hexagon_elementwise_add_kernel(
                     input2,
                     output,
                 ],
-                _SUPER_TARGET,
+                get_hexagon_target("v69"),
                 name=_PRIMFUNC_NAME,
             )
 

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,10 +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)
-                # 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)
+                    sch.mod,
+                    target=get_hexagon_target("v69"),
                 )
 
                 # Save a local copy of the Hexagon object code (in the form of a .so file)

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",
     )
 

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,
         )

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,
         )

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)