From f077faaaf6e8e1cb93b8afd2fab936effe04d5f4 Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Wed, 26 Oct 2022 08:45:57 -0700 Subject: [PATCH] [Hexagon]: Add upload function to hexagon session (#13161) * [Hexagon]: Add upload to hexagon session * lint * fix typo * fix serial device skips * fix test on device * add serial device to rpc key * update error name * fix comment * move create_session to launcher * add is_simulator * lint * address Eric comment on Session object * rebase with main --- python/tvm/contrib/hexagon/build.py | 195 +++++------------- python/tvm/contrib/hexagon/meta_schedule.py | 2 +- python/tvm/contrib/hexagon/pytest_plugin.py | 15 +- python/tvm/contrib/hexagon/session.py | 50 +++-- python/tvm/contrib/hexagon/tools.py | 2 + src/runtime/hexagon/rpc/hexagon/rpc_server.cc | 14 ++ .../hexagon/rpc/simulator/rpc_server.cc | 14 ++ .../python/contrib/test_hexagon/README_RPC.md | 2 +- .../contrib/test_hexagon/benchmark_util.py | 6 +- .../metaschedule_e2e/test_resnet50_fp16.py | 3 +- .../metaschedule_e2e/test_resnet50_int8.py | 2 +- .../test_benchmark_elemwise_add.py | 107 +++++----- .../test_hexagon/test_benchmark_maxpool2d.py | 2 +- .../test_hexagon/test_meta_schedule.py | 18 +- .../contrib/test_hexagon/test_models.py | 2 +- .../test_software_pipeline_async.py | 3 +- .../test_hexagon/topi/test_cast_slice.py | 5 +- 17 files changed, 196 insertions(+), 246 deletions(-) diff --git a/python/tvm/contrib/hexagon/build.py b/python/tvm/contrib/hexagon/build.py index c0e6439d0357..bc3b065dd941 100644 --- a/python/tvm/contrib/hexagon/build.py +++ b/python/tvm/contrib/hexagon/build.py @@ -31,14 +31,14 @@ import tempfile from typing import Union -import tvm from tvm.contrib.hexagon.hexagon_profiler import HexagonProfiler from ..._ffi import libinfo from .session import Session - +from .tools import HEXAGON_SIMULATOR_NAME HEXAGON_RPC_LIB_DIR = os.environ.get("HEXAGON_RPC_LIB_DIR") ANDROID_BASH_FILE_NAME = "android_bash.sh" +HEXAGON_REMOTE_DEVICE_KEY = "hexagon-dev" def _check_call_verbose(cmd, **kwargs) -> None: @@ -103,14 +103,9 @@ class HexagonLauncherRPC(metaclass=abc.ABCMeta): The basic flow of interaction with the launcher is launcher = HexagonLauncher(...) launcher.start_server() - with launcher.start_session() as session: + with launcher.create_session() as session: # Do something with the session launcher.stop_server() - """ - - HEXAGON_REMOTE_DEVICE_KEY = "hexagon-dev" - - """Configure HexagonLauncherRPC. Parameters ---------- @@ -129,7 +124,9 @@ class HexagonLauncherRPC(metaclass=abc.ABCMeta): used. """ - def __init__(self, rpc_info: dict, workspace: Union[str, pathlib.Path] = None): + def __init__( + self, rpc_info: dict, workspace: Union[str, pathlib.Path] = None, serial_number: str = None + ): self._rpc_info = { "rpc_tracker_host": "0.0.0.0", "rpc_tracker_port": 9190, @@ -138,7 +135,7 @@ def __init__(self, rpc_info: dict, workspace: Union[str, pathlib.Path] = None): } self._rpc_info.update(rpc_info) self._workspace = self._create_workspace(workspace) - self._device_key = self.HEXAGON_REMOTE_DEVICE_KEY + self._serial_number = serial_number @abc.abstractmethod def start_server(self): @@ -205,138 +202,6 @@ def _create_workspace(self, workspace: Union[str, pathlib.Path]) -> pathlib.Path workspace = os.path.join(base_dir, _get_test_directory_name()) return self._create_remote_directory(workspace) - def upload(self, local_path: Union[str, pathlib.Path], remote_filename: str) -> pathlib.Path: - """Upload a local file to the remote workspace. - - Parameters - ---------- - local_path : str or pathlib.Path - Path to the local file to be copied. - remote_filename : str - Name of the file in the remote workspace. - - Returns - ------- - pathlib.Path : - Uploaded file remote path. - """ - assert self._workspace - remote_file_path = self._workspace / remote_filename - self._copy_to_remote(local_path, str(remote_file_path)) - return remote_file_path - - def start_session(self, session_name: str = "hexagon-rpc") -> Session: - """Connect to the RPC server. - - Parameters - ---------- - session_name : str - RPC session name. - - Returns - ------- - Session : - The session object. - """ - hexagon_remote_kw = { - "host": self._rpc_info["rpc_tracker_host"], - "port": self._rpc_info["rpc_tracker_port"], - "priority": 0, - "timeout": 0, - "key": self._device_key, - } - return Session(self, hexagon_remote_kw, session_name=session_name) - - def load_module(self, module: Union[str, pathlib.Path, tvm.runtime.Module], session: Session): - """Load TVM module. - - Parameters - ---------- - module : Union[str, pathlib.Path, tvm.runtime.Module] - - The module to load. If `module` is a - `tvm.runtime.Module`, it will be uploaded to the remote - session and loaded. - - If the object passed is a string or pathlib.Path, it must - be a full path in the remote system. - - session : Session - - Remote session. The session must be established (via __enter__) - prior to calling this function. - - Returns - ------- - TVMModule : - TVM module object. - - """ - return session.load_module(module) - - def get_graph_executor( - self, - graph_json: str, - module: Union[str, pathlib.Path, tvm.runtime.Module], - session: Session, - ): - """Create a local GraphModule which consumes a remote libmod. - - Parameters - ---------- - graph_json : str - The string with the graph JSON. - module : Union[str, pathlib.Path, tvm.runtime.Module] - - The module to load. If `module` is a - `tvm.runtime.Module`, it will be uploaded to the remote - session and loaded. - - If the object passed is a string or pathlib.Path, it must - be a full path in the remote system. - session : Session - Remote session. The session must be established (via __enter__) - prior to calling this function. - - Returns - ------- - GraphModule : - Runtime graph module that can be used to execute the graph. - """ - return session.get_graph_executor(graph_json, module) - - def get_graph_debug_executor( - self, - graph_json: str, - module: Union[str, pathlib.Path, tvm.runtime.Module], - session: Session, - dump_root: Union[str, pathlib.Path] = None, - ): - """Create a local GraphModuleDebug which consumes a remote libmod. - - Parameters - ---------- - graph_json : str - The string with the graph JSON. - module : Union[str, pathlib.Path, tvm.runtime.Module] - - The module to load. If `module` is a - `tvm.runtime.Module`, it will be uploaded to the remote - session and loaded. - - If the object passed is a string or pathlib.Path, it must - be a full path in the remote system. - session : Session - Remote session. The session must be established (via __enter__) - prior to calling this function. - - Returns - ------- - GraphModuleDebug : - Runtime debug graph module that can be used to debug the graph. - """ - return session.get_graph_debug_executor(graph_json, module, dump_root=dump_root) - @abc.abstractmethod def get_profile_output( self, @@ -360,6 +225,31 @@ def get_profile_output( """ ... + def create_session(self, session_name: str = "hexagon-rpc") -> Session: + """Create an RPC session. + + Parameters + ---------- + session_name : str + RPC session name. + + Returns + ------- + Session : + The session object. + """ + hexagon_session_kw = { + "remote_workspace": self._workspace, + "rpc_tracker": (self._rpc_info["rpc_tracker_host"], self._rpc_info["rpc_tracker_port"]), + "rpc_server_key": self._rpc_info["device_key"], + "serial_number": self._serial_number, + "session_name": session_name, + } + return Session(**hexagon_session_kw) + + def is_simulator(self): + return self._serial_number == HEXAGON_SIMULATOR_NAME + class HexagonLauncherAndroid(HexagonLauncherRPC): """Hexagon Launcher for Android.""" @@ -402,6 +292,8 @@ def __init__( if not rpc_info.get("workspace_base"): rpc_info["workspace_base"] = self.ANDROID_HEXAGON_TEST_BASE_DIR self._serial_number = serial_number + assert self._serial_number != "", "Android serial number is not set." + adb_socket = rpc_info["adb_server_socket"] if rpc_info["adb_server_socket"] else "tcp:5037" self._adb_device_sub_cmd = ["adb", "-L", adb_socket, "-s", self._serial_number] self.forwarded_ports_ = [] @@ -409,8 +301,9 @@ def __init__( self._clear_logcat = clear_logcat self._sysmon_profile = sysmon_profile self._sysmon_process = None + rpc_info["device_key"] = HEXAGON_REMOTE_DEVICE_KEY + "." + self._serial_number - super(HexagonLauncherAndroid, self).__init__(rpc_info, workspace) + super(HexagonLauncherAndroid, self).__init__(rpc_info, workspace, self._serial_number) def _copy_to_remote( self, local_path: Union[str, pathlib.Path], remote_path: Union[str, pathlib.Path] @@ -442,7 +335,9 @@ def _copy_binaries(self): "", str(self._rpc_info["rpc_tracker_port"]) ) if "" in line: - line = line.replace("", self._device_key) + line = line.replace( + "", self._rpc_info["device_key"] + ) if "" in line: line = line.replace( "", str(self._rpc_info["rpc_server_port"]) @@ -691,12 +586,13 @@ def __init__(self, rpc_info: dict, workspace: Union[str, pathlib.Path] = None): Parameters are same as for HexagonLauncherRPC. """ - super(HexagonLauncherSimulator, self).__init__(rpc_info, workspace) self._toolchain = os.environ.get("HEXAGON_TOOLCHAIN") if not self._toolchain: raise RuntimeError("Please set HEXAGON_TOOLCHAIN env variable") - self._serial_number = "simulator" + self._serial_number = HEXAGON_SIMULATOR_NAME + + super(HexagonLauncherSimulator, self).__init__(rpc_info, workspace, self._serial_number) def _copy_to_remote( self, local_path: Union[str, pathlib.Path], remote_path: Union[str, pathlib.Path] @@ -740,18 +636,19 @@ def start_server(self): self._copy_to_remote(lib_dir / item, self._workspace / item) # Copy libc++ from the toolchain to the workspace self._copy_libcxx(self._workspace) - self._device_key = self.HEXAGON_REMOTE_DEVICE_KEY + "." + str(os.getpid()) + self._rpc_info["device_key"] = HEXAGON_REMOTE_DEVICE_KEY + "." + str(os.getpid()) rpc_tracker_host = self._rpc_info["rpc_tracker_host"] rpc_tracker_port = self._rpc_info["rpc_tracker_port"] rpc_server_port = self._rpc_info["rpc_server_port"] + device_key = self._rpc_info["device_key"] server_exe = os.path.join(".", "tvm_rpc_x86") args = [ "server", f"--tracker={rpc_tracker_host}:{rpc_tracker_port}", f"--port={rpc_server_port}", - f"--key={self._device_key}", + f"--key={device_key}", "--timeout=0", ] @@ -823,7 +720,7 @@ def HexagonLauncher( sysmon_profile: bool = False, ): """Creates a HexagonLauncher""" - if serial_number == "simulator": + if serial_number == HEXAGON_SIMULATOR_NAME: return HexagonLauncherSimulator(rpc_info, workspace) return HexagonLauncherAndroid( serial_number, rpc_info, workspace, hexagon_debug, clear_logcat, sysmon_profile diff --git a/python/tvm/contrib/hexagon/meta_schedule.py b/python/tvm/contrib/hexagon/meta_schedule.py index 8a4de74b6131..aaf3f8c7f8d5 100644 --- a/python/tvm/contrib/hexagon/meta_schedule.py +++ b/python/tvm/contrib/hexagon/meta_schedule.py @@ -100,7 +100,7 @@ def run(self, runner_inputs: List[RunnerInput]) -> List[RunnerFuture]: def _worker_func(hexagon_launcher, evaluator_config, alloc_repeat, artifact_path, args_info): - with hexagon_launcher.start_session() as session: + with hexagon_launcher.create_session() as session: device = session.device _, remote_path = os.path.split(artifact_path) uploaded = session.upload(artifact_path, remote_path) diff --git a/python/tvm/contrib/hexagon/pytest_plugin.py b/python/tvm/contrib/hexagon/pytest_plugin.py index 7ee16f50eab4..d462a65ef930 100644 --- a/python/tvm/contrib/hexagon/pytest_plugin.py +++ b/python/tvm/contrib/hexagon/pytest_plugin.py @@ -29,6 +29,7 @@ import tvm.rpc.tracker from tvm.contrib.hexagon.build import HexagonLauncher, HexagonLauncherRPC from tvm.contrib.hexagon.session import Session +from tvm.contrib.hexagon.tools import HEXAGON_SIMULATOR_NAME HEXAGON_TOOLCHAIN = "HEXAGON_TOOLCHAIN" TVM_TRACKER_HOST = "TVM_TRACKER_HOST" @@ -173,7 +174,7 @@ def hexagon_server_process( if android_serial_num is None: pytest.skip("ANDROID_SERIAL_NUMBER is not set.") - if android_serial_num == ["simulator"]: + if android_serial_num == [HEXAGON_SIMULATOR_NAME]: yield None else: # Requesting these fixtures sets up a local tracker, if one @@ -220,7 +221,7 @@ def pytest_configure(config): def pytest_configure_node(node): - # the master for each node fills slaveinput dictionary + # the master for each node fills node input dictionary # which pytest-xdist will transfer to the subprocess if node.config.iplist is not None: node.workerinput["device_adr"] = node.config.iplist.pop() @@ -240,7 +241,7 @@ def hexagon_launcher( """Initials and returns hexagon launcher which reuses RPC info and Android serial number.""" android_serial_num = android_serial_number() - if android_serial_num != ["simulator"]: + if android_serial_num != [HEXAGON_SIMULATOR_NAME]: rpc_info = hexagon_server_process["launcher"]._rpc_info else: rpc_info = { @@ -250,7 +251,7 @@ def hexagon_launcher( "adb_server_socket": adb_server_socket, } try: - if android_serial_num == ["simulator"]: + if android_serial_num == [HEXAGON_SIMULATOR_NAME]: launcher = HexagonLauncher(serial_number=android_serial_num[0], rpc_info=rpc_info) launcher.start_server() else: @@ -263,7 +264,7 @@ def hexagon_launcher( ) yield launcher finally: - if android_serial_num == ["simulator"]: + if android_serial_num == [HEXAGON_SIMULATOR_NAME]: launcher.stop_server() elif not hexagon_debug: launcher.cleanup_directory() @@ -274,7 +275,7 @@ def hexagon_session(hexagon_launcher: HexagonLauncherRPC) -> Session: if hexagon_launcher is None: yield None else: - with hexagon_launcher.start_session() as session: + with hexagon_launcher.create_session() as session: yield session @@ -289,7 +290,7 @@ def terminate_rpc_servers(): # yield happens every time. serial = os.environ.get(ANDROID_SERIAL_NUMBER) yield [] - if serial == ["simulator"]: + if serial == [HEXAGON_SIMULATOR_NAME]: os.system("ps ax | grep tvm_rpc_x86 | awk '{print $1}' | xargs kill") diff --git a/python/tvm/contrib/hexagon/session.py b/python/tvm/contrib/hexagon/session.py index d6ea51b53e17..466103f6e2c9 100644 --- a/python/tvm/contrib/hexagon/session.py +++ b/python/tvm/contrib/hexagon/session.py @@ -30,7 +30,7 @@ AOTExecutorFactoryModule, GraphExecutorFactoryModule, ) -from .tools import export_module +from .tools import export_module, HEXAGON_SIMULATOR_NAME class Session: @@ -38,11 +38,17 @@ class Session: Parameters ---------- - launcher : HexagonLauncherRPC - The launcher from which this session was started. + remote_workspace : Union[str, pathlib.Path] + Remote workspace path - remote_kw : dict - Remote configs for RPC tracker. + rpc_tracker : tuple(str, int) + RPC tracker host and port number. + + rpc_server_key : str + RPC server key on remote device. + + serial_number : str + Device serial number. `simulator` used for hexagon simulator. session_name : str Hexagon RPC session name. @@ -50,21 +56,28 @@ class Session: remote_stack_size_bytes : int The stack size of the remote device, to be passed to tvm.contrib.hexagon.create_hexagon_session. + + rpc_receive_buffer_size_bytes : int + RPC receive buffer size in bytes. """ def __init__( self, - launcher: "HexagonLauncherRPC", - remote_kw: dict, + remote_workspace: Union[str, pathlib.Path], + rpc_tracker: tuple, + rpc_server_key: str, + serial_number: str, session_name: str = "hexagon-rpc", remote_stack_size_bytes: int = 256 * 1024, # Min size for main thread in QuRT/sim rpc_receive_buffer_size_bytes: int = 256 * 1024 * 1024, # Size for passing hexagon tests ): - self._launcher = launcher + self._workspace = str(remote_workspace) + self._rpc_tracker = rpc_tracker + self._rpc_server_key = rpc_server_key + self._serial_number = serial_number self._session_name: str = session_name self._remote_stack_size_bytes: int = remote_stack_size_bytes self._rpc_receive_buffer_size_bytes: int = rpc_receive_buffer_size_bytes - self._remote_kw: dict = remote_kw self._rpc = None self._requires_cpu_device = False self._device = None @@ -74,12 +87,12 @@ def __enter__(self): # Already initialized return self - tracker = _rpc.connect_tracker(self._remote_kw["host"], self._remote_kw["port"]) + tracker = _rpc.connect_tracker(self._rpc_tracker[0], self._rpc_tracker[1]) try: self._rpc = tracker.request( - self._remote_kw["key"], - priority=self._remote_kw["priority"], - session_timeout=self._remote_kw["timeout"], + self._rpc_server_key, + priority=0, + session_timeout=0, session_constructor_args=[ "tvm.contrib.hexagon.create_hexagon_session", self._session_name, @@ -124,6 +137,9 @@ def device(self): return self._device + def is_simulator(self): + return self._serial_number == HEXAGON_SIMULATOR_NAME + def get_function(self, name): return self._rpc.get_function(name) @@ -142,7 +158,12 @@ def upload(self, local_path: Union[str, pathlib.Path], remote_filename: str) -> pathlib.Path : Uploaded file remote path. """ - return self._launcher.upload(local_path, remote_filename) + upload_func = self._rpc.get_function("tvm.rpc.server.upload") + remote_path = f"{self._workspace}/{remote_filename}" + with open(local_path, mode="rb") as src_f: + data = bytearray(src_f.read()) + upload_func(remote_path, data) + return remote_path def load_module(self, module: Union[str, pathlib.Path, tvm.runtime.Module]): """Load TVM module. @@ -206,7 +227,6 @@ def get_graph_executor( Runtime graph module that can be used to execute the graph. """ - graph_mod = self.load_module(module_name) self._set_device_type(graph_mod) return tvm.contrib.graph_executor.create(graph_json, graph_mod, self.device) diff --git a/python/tvm/contrib/hexagon/tools.py b/python/tvm/contrib/hexagon/tools.py index 3f4adb90f645..8c37261744d5 100644 --- a/python/tvm/contrib/hexagon/tools.py +++ b/python/tvm/contrib/hexagon/tools.py @@ -53,6 +53,8 @@ pathlib.Path(HEXAGON_SDK_ROOT) / "incs" / "stddef", ] +HEXAGON_SIMULATOR_NAME = "simulator" + def register_linker(f): """Register a function that will return the path to the Hexagon linker.""" diff --git a/src/runtime/hexagon/rpc/hexagon/rpc_server.cc b/src/runtime/hexagon/rpc/hexagon/rpc_server.cc index 41c63d0affeb..f39944615bfd 100644 --- a/src/runtime/hexagon/rpc/hexagon/rpc_server.cc +++ b/src/runtime/hexagon/rpc/hexagon/rpc_server.cc @@ -32,6 +32,7 @@ extern "C" { #include #include +#include #include #include @@ -342,3 +343,16 @@ TVM_REGISTER_GLOBAL("tvm.hexagon.get_profile_output") *rv = false; } }); + +void SaveBinaryToFile(const std::string& file_name, const std::string& data) { + std::ofstream fs(file_name, std::ios::out | std::ios::binary); + ICHECK(!fs.fail()) << "Cannot open " << file_name; + fs.write(&data[0], data.length()); +} + +TVM_REGISTER_GLOBAL("tvm.rpc.server.upload") + .set_body([](tvm::runtime::TVMArgs args, tvm::runtime::TVMRetValue* rv) { + std::string file_name = args[0]; + std::string data = args[1]; + SaveBinaryToFile(file_name, data); + }); diff --git a/src/runtime/hexagon/rpc/simulator/rpc_server.cc b/src/runtime/hexagon/rpc/simulator/rpc_server.cc index 41bb2da6f8b1..f4370bd3c88c 100644 --- a/src/runtime/hexagon/rpc/simulator/rpc_server.cc +++ b/src/runtime/hexagon/rpc/simulator/rpc_server.cc @@ -23,6 +23,7 @@ #include #include #include +#include #include #include @@ -349,3 +350,16 @@ TVM_REGISTER_GLOBAL("tvm.hexagon.get_profile_output") *rv = false; } }); + +void SaveBinaryToFile(const std::string& file_name, const std::string& data) { + std::ofstream fs(file_name, std::ios::out | std::ios::binary); + ICHECK(!fs.fail()) << "Cannot open " << file_name; + fs.write(&data[0], data.length()); +} + +TVM_REGISTER_GLOBAL("tvm.rpc.server.upload") + .set_body([](tvm::runtime::TVMArgs args, tvm::runtime::TVMRetValue* rv) { + std::string file_name = args[0]; + std::string data = args[1]; + SaveBinaryToFile(file_name, data); + }); diff --git a/tests/python/contrib/test_hexagon/README_RPC.md b/tests/python/contrib/test_hexagon/README_RPC.md index 348be2d9e457..228922d3f184 100644 --- a/tests/python/contrib/test_hexagon/README_RPC.md +++ b/tests/python/contrib/test_hexagon/README_RPC.md @@ -60,7 +60,7 @@ subprocess.Popen( ./tvm_rpc_android server --port= --tracker=: --key=& ``` -When we do `launcher.start_session()` , a remote RPC session between x86 and android is established via this line: +When we do `launcher.create_session()` , a remote RPC session between x86 and android is established via this line: [https://github.com/apache/tvm/blob/0c0245ae2230fa07d3e4b8be490fc9c88965730c/python/tvm/contrib/hexagon/session.py#L57-L67](https://github.com/apache/tvm/blob/0c0245ae2230fa07d3e4b8be490fc9c88965730c/python/tvm/contrib/hexagon/session.py#L57-L67) diff --git a/tests/python/contrib/test_hexagon/benchmark_util.py b/tests/python/contrib/test_hexagon/benchmark_util.py index 0ded60dc498b..0ccbe514326c 100644 --- a/tests/python/contrib/test_hexagon/benchmark_util.py +++ b/tests/python/contrib/test_hexagon/benchmark_util.py @@ -23,8 +23,10 @@ import pytest +from tvm.contrib.hexagon.tools import HEXAGON_SIMULATOR_NAME -def skip_bencharks_flag_and_reason(): + +def skip_benchmarks_flag_and_reason(): """ Returns one of these tuples: (False, '') or @@ -37,7 +39,7 @@ def skip_bencharks_flag_and_reason(): """ asn = os.environ.get("ANDROID_SERIAL_NUMBER") - if asn == "simulator": + if asn == HEXAGON_SIMULATOR_NAME: return (True, "Skipping benchmarks when ANDROID_SERIAL_NUMBER='simluator'") return (False, "") diff --git a/tests/python/contrib/test_hexagon/metaschedule_e2e/test_resnet50_fp16.py b/tests/python/contrib/test_hexagon/metaschedule_e2e/test_resnet50_fp16.py index 4fe21c564330..84a33b9c80d3 100644 --- a/tests/python/contrib/test_hexagon/metaschedule_e2e/test_resnet50_fp16.py +++ b/tests/python/contrib/test_hexagon/metaschedule_e2e/test_resnet50_fp16.py @@ -25,6 +25,7 @@ from tvm import meta_schedule as ms from tvm.contrib.hexagon.meta_schedule import get_hexagon_local_builder, get_hexagon_rpc_runner from tvm.relay.backend import Executor + from ..infrastructure import get_hexagon_target @@ -103,7 +104,7 @@ def test_resnet50(hexagon_launcher): llvm_graph_mod.run() ref_result = llvm_graph_mod.get_output(0).numpy() - with hexagon_launcher.start_session() as session: + with hexagon_launcher.create_session() as session: graph_mod = session.get_executor_from_factory(hexagon_lowered) graph_mod.set_input(input_name, inp.copy()) diff --git a/tests/python/contrib/test_hexagon/metaschedule_e2e/test_resnet50_int8.py b/tests/python/contrib/test_hexagon/metaschedule_e2e/test_resnet50_int8.py index 4c8d91dd27ef..a541c25f3cbc 100644 --- a/tests/python/contrib/test_hexagon/metaschedule_e2e/test_resnet50_int8.py +++ b/tests/python/contrib/test_hexagon/metaschedule_e2e/test_resnet50_int8.py @@ -163,7 +163,7 @@ def test_resnet50(hexagon_launcher): params=params, ) - with hexagon_launcher.start_session() as session: + with hexagon_launcher.create_session() as session: graph_mod = session.get_executor_from_factory(hexagon_lowered) graph_mod.set_input(input_name, inp.copy()) graph_mod.run() 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 3dcb9a880e00..ee58514569ae 100644 --- a/tests/python/contrib/test_hexagon/test_benchmark_elemwise_add.py +++ b/tests/python/contrib/test_hexagon/test_benchmark_elemwise_add.py @@ -26,13 +26,13 @@ import tvm.script import tvm.testing -from tvm.contrib.hexagon.build import HexagonLauncherRPC +from tvm.contrib.hexagon.session import Session 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() +_SHOULD_SKIP_BENCHMARKS, _SKIP_BENCHMARKS_REASON = bu.skip_benchmarks_flag_and_reason() # This is a fixed detail of the v68 architecture. HVX_VECTOR_BYTES = 128 @@ -160,7 +160,7 @@ def main(a: T.handle, b: T.handle, c: T.handle): def _benchmark_hexagon_elementwise_add_kernel( - hexagon_launcher: HexagonLauncherRPC, shape: list, dtype: str, mem_scope: str + hexagon_session: Session, shape: list, dtype: str, mem_scope: str ): """ Generate and benchmark a single elementwise-add kernel for Hexagon. @@ -230,7 +230,7 @@ def _benchmark_hexagon_elementwise_add_kernel( # Upload the .so to the Android device's file system (or wherever is appropriate # when using the Hexagon simulator)... target_dso_binary_filename = "test_binary.so" - target_dso_binary_pathname = hexagon_launcher.upload( + target_dso_binary_pathname = hexagon_session.upload( host_dso_binary_path, target_dso_binary_filename ) @@ -241,58 +241,57 @@ def _benchmark_hexagon_elementwise_add_kernel( host_numpy_output_data_expected, ) = _get_elemwise_add_reference_value_tensors(shape, dtype) - with hexagon_launcher.start_session() as sess: - # On the target device / simulator, make our Hexagon-native shared object - # available for use... - loaded_hexagon_module: tvm.runtime.module.Module = hexagon_launcher.load_module( - target_dso_binary_pathname, sess - ) + # On the target device / simulator, make our Hexagon-native shared object + # available for use... + loaded_hexagon_module: tvm.runtime.module.Module = hexagon_session.load_module( + target_dso_binary_pathname + ) - # Create the target-side tensors to hold the primfunc's inputs and outputs... - input1_data = tvm.nd.empty(shape, dtype, sess.device, mem_scope) - input2_data = tvm.nd.empty(shape, dtype, sess.device, mem_scope) - output_data = tvm.nd.empty(shape, dtype, sess.device, mem_scope) + # Create the target-side tensors to hold the primfunc's inputs and outputs... + input1_data = tvm.nd.empty(shape, dtype, hexagon_session.device, mem_scope) + input2_data = tvm.nd.empty(shape, dtype, hexagon_session.device, mem_scope) + output_data = tvm.nd.empty(shape, dtype, hexagon_session.device, mem_scope) - # Populate the primfunc's input tensors... - input1_data.copyfrom(host_numpy_input1_data) - input2_data.copyfrom(host_numpy_input2_data) + # Populate the primfunc's input tensors... + input1_data.copyfrom(host_numpy_input1_data) + input2_data.copyfrom(host_numpy_input2_data) - # Actually benchmark the primfunc... - timer = loaded_hexagon_module.time_evaluator( - "main", sess.device, number=10, repeat=1 + # Actually benchmark the primfunc... + timer = loaded_hexagon_module.time_evaluator( + "main", hexagon_session.device, number=10, repeat=1 + ) + timing_result = timer(input1_data, input2_data, output_data) + + print(f"TIMING RESULT: {timing_result}") + log_file.write(f"TIMING RESULT: {timing_result}\n") + + # Verify that the computation actually happened, and produced the correct result. + result = output_data.numpy() + + if dtype == "float16": + # These are the closest tolerance we currently expect / require for these + # kernels. They may be changed in the future. + rel_tolerance = 0.005 + abs_tolerance = 2.0 + elif dtype == "int8": + rel_tolerance = 0 + abs_tolerance = 0 + else: + raise Exception(f"Unexpected dtype: {dtype}") + + # TODO: We're assuming that *any* assertion thrown by 'assert_allclose' is because + # the numerical differences were too large. But ideally this code would + # differentiate between (a) numerical difference errors, which should simply be + # recorded as a failed benchmark run, vs. (b) more serious errors that should + # kill the overall script. + try: + tvm.testing.assert_allclose( + result, host_numpy_output_data_expected, rel_tolerance, abs_tolerance ) - timing_result = timer(input1_data, input2_data, output_data) - - print(f"TIMING RESULT: {timing_result}") - log_file.write(f"TIMING RESULT: {timing_result}\n") - - # Verify that the computation actually happened, and produced the correct result. - result = output_data.numpy() - - if dtype == "float16": - # These are the closest tolerance we currently expect / require for these - # kernels. They may be changed in the future. - rel_tolerance = 0.005 - abs_tolerance = 2.0 - elif dtype == "int8": - rel_tolerance = 0 - abs_tolerance = 0 - else: - raise Exception(f"Unexpected dtype: {dtype}") - - # TODO: We're assuming that *any* assertion thrown by 'assert_allclose' is because - # the numerical differences were too large. But ideally this code would - # differentiate between (a) numerical difference errors, which should simply be - # recorded as a failed benchmark run, vs. (b) more serious errors that should - # kill the overall script. - try: - tvm.testing.assert_allclose( - result, host_numpy_output_data_expected, rel_tolerance, abs_tolerance - ) - except AssertionError as err: - raise bu.NumericalAccuracyException(str(err)) - - _BT.record_success(timing_result, **keys_dict) + except AssertionError as err: + raise bu.NumericalAccuracyException(str(err)) + + _BT.record_success(timing_result, **keys_dict) except bu.NumericalAccuracyException as err: print() @@ -377,7 +376,7 @@ def _get_elemwise_add_reference_value_tensors(shape: list, dtype: str): @pytest.mark.skipif(_SHOULD_SKIP_BENCHMARKS, reason=_SKIP_BENCHMARKS_REASON) @tvm.testing.requires_hexagon -def test_elemwise_add(hexagon_launcher: HexagonLauncherRPC): +def test_elemwise_add(hexagon_session: Session): """Main elementwise add test function""" for dtype in [ "int8", @@ -411,7 +410,7 @@ def test_elemwise_add(hexagon_launcher: HexagonLauncherRPC): ] print() - _benchmark_hexagon_elementwise_add_kernel(hexagon_launcher, shape, dtype, mem_scope) + _benchmark_hexagon_elementwise_add_kernel(hexagon_session, shape, dtype, mem_scope) print("-" * 80) print(f"OUTPUT DIRECTORY: {_HOST_OUTPUT_DIR}") diff --git a/tests/python/contrib/test_hexagon/test_benchmark_maxpool2d.py b/tests/python/contrib/test_hexagon/test_benchmark_maxpool2d.py index a22b85ee42a2..7e8a6d79f492 100644 --- a/tests/python/contrib/test_hexagon/test_benchmark_maxpool2d.py +++ b/tests/python/contrib/test_hexagon/test_benchmark_maxpool2d.py @@ -60,7 +60,7 @@ # E.g., it doesn't allow: @pytest.mark.usefixtures("bu.benchmark_group") benchmark_group = bu.benchmark_group -_SHOULD_SKIP_BENCHMARKS, _SKIP_BENCHMARKS_REASON = bu.skip_bencharks_flag_and_reason() +_SHOULD_SKIP_BENCHMARKS, _SKIP_BENCHMARKS_REASON = bu.skip_benchmarks_flag_and_reason() def _ceil_div(numerator, denominator): diff --git a/tests/python/contrib/test_hexagon/test_meta_schedule.py b/tests/python/contrib/test_hexagon/test_meta_schedule.py index 6e12f4b205d1..a7f4cbc39cb1 100644 --- a/tests/python/contrib/test_hexagon/test_meta_schedule.py +++ b/tests/python/contrib/test_hexagon/test_meta_schedule.py @@ -62,7 +62,7 @@ def main( # type: ignore # pylint: disable=no-self-argument @tvm.testing.requires_hexagon def test_builder_runner(hexagon_launcher): - if hexagon_launcher._serial_number == "simulator": + if hexagon_launcher.is_simulator(): pytest.skip(msg="Tuning on simulator not supported.") mod = MatmulModule @@ -175,7 +175,7 @@ def verify_dense(sch, target, M, N, K, hexagon_session): @tvm.testing.requires_hexagon def test_vrmpy_dense(hexagon_launcher): - if hexagon_launcher._serial_number == "simulator": + if hexagon_launcher.is_simulator(): pytest.skip(msg="Tuning on simulator not supported.") do_tune = True @@ -213,7 +213,7 @@ def schedule_dense_for_tune(sch): ) sch = ms.tir_integration.compile_tir(database, workload, target) - with hexagon_launcher.start_session() as session: + with hexagon_launcher.create_session() as session: verify_dense(sch, get_hexagon_target("v68"), M, N, K, session) @@ -273,7 +273,7 @@ def main( # type: ignore @tvm.testing.requires_hexagon def test_vrmpy_dense_auto_tensorize(hexagon_launcher): - if hexagon_launcher._serial_number == "simulator": + if hexagon_launcher.is_simulator(): pytest.skip(msg="Tuning on simulator not supported.") M, N, K = 128, 768, 768 @@ -330,13 +330,13 @@ def test_vrmpy_dense_auto_tensorize(hexagon_launcher): else: sch = tvm.tir.Schedule(Module_vrmpy_auto_tensorize, debug_mask="all") - with hexagon_launcher.start_session() as session: + with hexagon_launcher.create_session() as session: verify_dense(sch, get_hexagon_target("v68"), M, N, K, session) @tvm.testing.requires_hexagon def test_conv2d_relay_auto_schedule(hexagon_launcher): - if hexagon_launcher._serial_number == "simulator": + if hexagon_launcher.is_simulator(): pytest.skip(msg="Tuning on simulator not supported.") I, O, H, W = 64, 64, 56, 56 @@ -397,7 +397,7 @@ def test_conv2d_relay_auto_schedule(hexagon_launcher): target=target, ) - with hexagon_launcher.start_session() as session: + with hexagon_launcher.create_session() as session: rt_mod = session.get_executor_from_factory(lib) rt_mod.set_input("data", data_np) @@ -416,7 +416,7 @@ def test_dense_relay_auto_schedule(hexagon_launcher): This is for testing RewriteLayout postproc. Without this postproc, dense on Hexagon is extremely slow. """ - if hexagon_launcher._serial_number == "simulator": + if hexagon_launcher.is_simulator(): pytest.skip(msg="Tuning on simulator not supported.") target_hexagon = tvm.target.hexagon("v69") @@ -456,7 +456,7 @@ def test_dense_relay_auto_schedule(hexagon_launcher): target=target, ) - with hexagon_launcher.start_session() as session: + with hexagon_launcher.create_session() as session: rt_mod = session.get_executor_from_factory(lib) rt_mod.set_input("data", data_np) diff --git a/tests/python/contrib/test_hexagon/test_models.py b/tests/python/contrib/test_hexagon/test_models.py index f4495f849fab..db578d1057a6 100644 --- a/tests/python/contrib/test_hexagon/test_models.py +++ b/tests/python/contrib/test_hexagon/test_models.py @@ -90,7 +90,7 @@ def test_mobilenet(hexagon_session: Session): @tvm.testing.requires_hexagon def test_mobilenet_aot(hexagon_session: Session, aot_host_target, aot_target, enable_usmp): """Test mobilenet with aot executor""" - if hexagon_session._launcher._serial_number == "simulator": + if hexagon_session.is_simulator(): pytest.skip(msg="Skip on simulator due to long runtime.") dtype = "float32" 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 a883a9a251e3..f80a579f58fe 100644 --- a/tests/python/contrib/test_hexagon/test_software_pipeline_async.py +++ b/tests/python/contrib/test_hexagon/test_software_pipeline_async.py @@ -21,7 +21,6 @@ import tvm from tvm import tir -from tvm.contrib.hexagon.session import Session from tvm.script import tir as T from .infrastructure import get_hexagon_target @@ -181,7 +180,7 @@ def test_async_software_pipeline(hexagon_launcher, comp_type, data, reference, s # tvm.lower(schedule.mod["main"]).show() func = tvm.build(schedule.mod["main"], target=get_hexagon_target("v68")) - with hexagon_launcher.start_session() as hexagon_session: + with hexagon_launcher.create_session() as hexagon_session: dev = hexagon_session.device mod = hexagon_session.load_module(func) out = tvm.nd.array(out_np, device=dev) 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 326370eb72d7..7f59e3ffa7fd 100644 --- a/tests/python/contrib/test_hexagon/topi/test_cast_slice.py +++ b/tests/python/contrib/test_hexagon/topi/test_cast_slice.py @@ -22,6 +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, get_hexagon_target @@ -74,7 +75,7 @@ def test_cast_fp16_fp32_slice( """ Top level testing function for cast fp16 to fp32 """ - if hexagon_session._launcher._serial_number != "simulator": + if hexagon_session.is_simulator(): pytest.skip(msg="Due to https://github.com/apache/tvm/issues/11957") cast_input = te.placeholder(input_shape, name="A", dtype=dtype) @@ -160,7 +161,7 @@ def test_cast_fp32_fp16_slice( """ Top level testing function for cast fp32 to fp16 """ - if hexagon_session._launcher._serial_number != "simulator": + if hexagon_session.is_simulator(): pytest.skip(msg="Due to https://github.com/apache/tvm/issues/11957") cast_input = te.placeholder(input_shape, name="A", dtype=dtype)