From 057d06fe934414bbb8c79ec6b103c1c1c040605c Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Mon, 17 Oct 2022 23:43:55 +0000 Subject: [PATCH 01/13] [Hexagon]: Add upload to hexagon session --- python/tvm/contrib/hexagon/build.py | 149 ++---------------- python/tvm/contrib/hexagon/meta_schedule.py | 3 +- python/tvm/contrib/hexagon/pytest_plugin.py | 4 +- python/tvm/contrib/hexagon/session.py | 16 +- src/runtime/hexagon/rpc/hexagon/rpc_server.cc | 12 ++ .../hexagon/rpc/simulator/rpc_server.cc | 12 ++ .../python/contrib/test_hexagon/README_RPC.md | 2 +- .../contrib/test_hexagon/benchmark_util.py | 2 +- .../metaschedule_e2e/test_resnet50_fp16.py | 4 +- .../metaschedule_e2e/test_resnet50_int8.py | 2 +- .../test_benchmark_elemwise_add.py | 107 +++++++------ .../test_hexagon/test_meta_schedule.py | 20 +-- .../contrib/test_hexagon/test_models.py | 3 +- .../test_software_pipeline_async.py | 4 +- .../test_hexagon/topi/test_cast_slice.py | 6 +- 15 files changed, 124 insertions(+), 222 deletions(-) diff --git a/python/tvm/contrib/hexagon/build.py b/python/tvm/contrib/hexagon/build.py index c0e6439d0357..348463c962e6 100644 --- a/python/tvm/contrib/hexagon/build.py +++ b/python/tvm/contrib/hexagon/build.py @@ -34,11 +34,11 @@ import tvm from tvm.contrib.hexagon.hexagon_profiler import HexagonProfiler from ..._ffi import libinfo -from .session import Session 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,13 +103,11 @@ 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 session.create_session() as session: # Do something with the session launcher.stop_server() """ - HEXAGON_REMOTE_DEVICE_KEY = "hexagon-dev" - """Configure HexagonLauncherRPC. Parameters @@ -138,7 +136,6 @@ 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 @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, @@ -409,6 +274,7 @@ def __init__( self._clear_logcat = clear_logcat self._sysmon_profile = sysmon_profile self._sysmon_process = None + rpc_info["device_key"] = HEXAGON_REMOTE_DEVICE_KEY super(HexagonLauncherAndroid, self).__init__(rpc_info, workspace) @@ -442,7 +308,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"]) @@ -740,18 +608,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", ] diff --git a/python/tvm/contrib/hexagon/meta_schedule.py b/python/tvm/contrib/hexagon/meta_schedule.py index 8a4de74b6131..7ea2c9c272a7 100644 --- a/python/tvm/contrib/hexagon/meta_schedule.py +++ b/python/tvm/contrib/hexagon/meta_schedule.py @@ -35,6 +35,7 @@ from .build import HexagonLauncherRPC from .tools import export_module +from .session import create_session @derived_object @@ -100,7 +101,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 create_session(hexagon_launcher._workspace, hexagon_launcher._rpc_info) 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..6b8bb61e8b0d 100644 --- a/python/tvm/contrib/hexagon/pytest_plugin.py +++ b/python/tvm/contrib/hexagon/pytest_plugin.py @@ -28,7 +28,7 @@ import tvm import tvm.rpc.tracker from tvm.contrib.hexagon.build import HexagonLauncher, HexagonLauncherRPC -from tvm.contrib.hexagon.session import Session +from tvm.contrib.hexagon.session import Session, create_session HEXAGON_TOOLCHAIN = "HEXAGON_TOOLCHAIN" TVM_TRACKER_HOST = "TVM_TRACKER_HOST" @@ -274,7 +274,7 @@ def hexagon_session(hexagon_launcher: HexagonLauncherRPC) -> Session: if hexagon_launcher is None: yield None else: - with hexagon_launcher.start_session() as session: + with create_session(hexagon_launcher._workspace, hexagon_launcher._rpc_info) as session: yield session diff --git a/python/tvm/contrib/hexagon/session.py b/python/tvm/contrib/hexagon/session.py index d6ea51b53e17..c10fe83d19df 100644 --- a/python/tvm/contrib/hexagon/session.py +++ b/python/tvm/contrib/hexagon/session.py @@ -31,6 +31,7 @@ GraphExecutorFactoryModule, ) from .tools import export_module +from .build import HEXAGON_REMOTE_DEVICE_KEY class Session: @@ -38,9 +39,6 @@ class Session: Parameters ---------- - launcher : HexagonLauncherRPC - The launcher from which this session was started. - remote_kw : dict Remote configs for RPC tracker. @@ -54,13 +52,13 @@ class Session: def __init__( self, - launcher: "HexagonLauncherRPC", + remote_workspace: Union[str, pathlib.Path], remote_kw: dict, 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._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 @@ -142,7 +140,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 +209,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/src/runtime/hexagon/rpc/hexagon/rpc_server.cc b/src/runtime/hexagon/rpc/hexagon/rpc_server.cc index 41c63d0affeb..77c3b8e94d40 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 @@ -341,4 +342,15 @@ TVM_REGISTER_GLOBAL("tvm.hexagon.get_profile_output") HEXAGON_PRINT(ERROR, "ERROR: Unsupported profiling mode: %s", profiling_mode.c_str()); *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..2dd36ee287db 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 @@ -348,4 +349,15 @@ TVM_REGISTER_GLOBAL("tvm.hexagon.get_profile_output") HEXAGON_PRINT(ERROR, "ERROR: Unsupported profiling mode: %s", profiling_mode.c_str()); *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..4aff7c203e70 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 `session.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..8ffdfe5090ca 100644 --- a/tests/python/contrib/test_hexagon/benchmark_util.py +++ b/tests/python/contrib/test_hexagon/benchmark_util.py @@ -24,7 +24,7 @@ import pytest -def skip_bencharks_flag_and_reason(): +def skip_benchmarks_flag_and_reason(): """ Returns one of these tuples: (False, '') or 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..2f16d6050bf4 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,8 @@ 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 tvm.contrib.hexagon.session import create_session + from ..infrastructure import get_hexagon_target @@ -103,7 +105,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 create_session(hexagon_launcher._workspace, hexagon_launcher._rpc_info) 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..8fac962b9b53 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 create_session(hexagon_launcher._workspace, hexagon_launcher._rpc_info) 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..4f172ec52aff 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 python.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, hexagon_session + ) - # 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_meta_schedule.py b/tests/python/contrib/test_hexagon/test_meta_schedule.py index 6e12f4b205d1..5105bf2749c8 100644 --- a/tests/python/contrib/test_hexagon/test_meta_schedule.py +++ b/tests/python/contrib/test_hexagon/test_meta_schedule.py @@ -35,6 +35,8 @@ from tvm.script import tir as T from tvm.tir import FloatImm from tvm.tir.tensor_intrin.hexagon import VRMPY_u8u8i32_INTRIN +from tvm.contrib.hexagon.pytest_plugin import android_serial_number +from tvm.contrib.hexagon.session import create_session from .infrastructure import get_hexagon_target @@ -62,7 +64,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 android_serial_number() == "simulator": pytest.skip(msg="Tuning on simulator not supported.") mod = MatmulModule @@ -175,7 +177,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 android_serial_number() == "simulator": pytest.skip(msg="Tuning on simulator not supported.") do_tune = True @@ -213,7 +215,7 @@ def schedule_dense_for_tune(sch): ) sch = ms.tir_integration.compile_tir(database, workload, target) - with hexagon_launcher.start_session() as session: + with create_session(hexagon_launcher._workspace, hexagon_launcher._rpc_info) as session: verify_dense(sch, get_hexagon_target("v68"), M, N, K, session) @@ -273,7 +275,7 @@ def main( # type: ignore @tvm.testing.requires_hexagon def test_vrmpy_dense_auto_tensorize(hexagon_launcher): - if hexagon_launcher._serial_number == "simulator": + if android_serial_number() == "simulator": pytest.skip(msg="Tuning on simulator not supported.") M, N, K = 128, 768, 768 @@ -330,13 +332,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 create_session(hexagon_launcher._workspace, hexagon_launcher._rpc_info) 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 android_serial_number() == "simulator": pytest.skip(msg="Tuning on simulator not supported.") I, O, H, W = 64, 64, 56, 56 @@ -397,7 +399,7 @@ def test_conv2d_relay_auto_schedule(hexagon_launcher): target=target, ) - with hexagon_launcher.start_session() as session: + with create_session(hexagon_launcher._workspace, hexagon_launcher._rpc_info) as session: rt_mod = session.get_executor_from_factory(lib) rt_mod.set_input("data", data_np) @@ -416,7 +418,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 android_serial_number() == "simulator": pytest.skip(msg="Tuning on simulator not supported.") target_hexagon = tvm.target.hexagon("v69") @@ -456,7 +458,7 @@ def test_dense_relay_auto_schedule(hexagon_launcher): target=target, ) - with hexagon_launcher.start_session() as session: + with create_session(hexagon_launcher._workspace, hexagon_launcher._rpc_info) 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..7d60072a0367 100644 --- a/tests/python/contrib/test_hexagon/test_models.py +++ b/tests/python/contrib/test_hexagon/test_models.py @@ -24,6 +24,7 @@ from tvm import relay from tvm.contrib.hexagon.session import Session from tvm.relay.backend import Executor, Runtime +from tvm.contrib.hexagon.pytest_plugin import android_serial_number from .infrastructure import get_hexagon_target @@ -90,7 +91,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 android_serial_number() == "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..348a3c799eb0 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,7 @@ import tvm from tvm import tir -from tvm.contrib.hexagon.session import Session +from tvm.contrib.hexagon.session import create_session from tvm.script import tir as T from .infrastructure import get_hexagon_target @@ -181,7 +181,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 create_session(hexagon_launcher._workspace, hexagon_launcher._rpc_info) 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..5e53da92c49e 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,8 @@ import tvm.testing from tvm import te import tvm.topi.hexagon.slice_ops as sl +from tvm.contrib.hexagon.pytest_plugin import android_serial_number + from ..infrastructure import allocate_hexagon_array, transform_numpy, get_hexagon_target @@ -74,7 +76,7 @@ def test_cast_fp16_fp32_slice( """ Top level testing function for cast fp16 to fp32 """ - if hexagon_session._launcher._serial_number != "simulator": + if android_serial_number() != "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 +162,7 @@ def test_cast_fp32_fp16_slice( """ Top level testing function for cast fp32 to fp16 """ - if hexagon_session._launcher._serial_number != "simulator": + if android_serial_number() != "simulator": pytest.skip(msg="Due to https://github.com/apache/tvm/issues/11957") cast_input = te.placeholder(input_shape, name="A", dtype=dtype) From 5f2a393b410c39284fee671b0df0e3e70b6964ea Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Thu, 20 Oct 2022 21:21:23 +0000 Subject: [PATCH 02/13] lint --- python/tvm/contrib/hexagon/build.py | 3 --- python/tvm/contrib/hexagon/session.py | 1 - 2 files changed, 4 deletions(-) diff --git a/python/tvm/contrib/hexagon/build.py b/python/tvm/contrib/hexagon/build.py index 348463c962e6..7582dcb7becc 100644 --- a/python/tvm/contrib/hexagon/build.py +++ b/python/tvm/contrib/hexagon/build.py @@ -106,9 +106,6 @@ class HexagonLauncherRPC(metaclass=abc.ABCMeta): with session.create_session() as session: # Do something with the session launcher.stop_server() - """ - - """Configure HexagonLauncherRPC. Parameters ---------- diff --git a/python/tvm/contrib/hexagon/session.py b/python/tvm/contrib/hexagon/session.py index c10fe83d19df..47dee1d29606 100644 --- a/python/tvm/contrib/hexagon/session.py +++ b/python/tvm/contrib/hexagon/session.py @@ -31,7 +31,6 @@ GraphExecutorFactoryModule, ) from .tools import export_module -from .build import HEXAGON_REMOTE_DEVICE_KEY class Session: From a4f7bffaa4ab433ab3b682d1d36b1d5abf98003f Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Thu, 20 Oct 2022 23:32:37 +0000 Subject: [PATCH 03/13] fix typo --- .../python/contrib/test_hexagon/test_benchmark_elemwise_add.py | 2 +- tests/python/contrib/test_hexagon/test_benchmark_maxpool2d.py | 2 +- 2 files changed, 2 insertions(+), 2 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 4f172ec52aff..769e9bbfce87 100644 --- a/tests/python/contrib/test_hexagon/test_benchmark_elemwise_add.py +++ b/tests/python/contrib/test_hexagon/test_benchmark_elemwise_add.py @@ -26,7 +26,7 @@ import tvm.script import tvm.testing -from python.tvm.contrib.hexagon.session import Session +from tvm.contrib.hexagon.session import Session from tvm.script import tir as T from . import benchmark_util as bu 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): From 0ce1d1e94eaa3368704fa2b520529805dd51bb26 Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Fri, 21 Oct 2022 06:03:49 +0000 Subject: [PATCH 04/13] fix serial device skips --- .../python/contrib/test_hexagon/test_meta_schedule.py | 10 +++++----- tests/python/contrib/test_hexagon/test_models.py | 2 +- .../contrib/test_hexagon/topi/test_cast_slice.py | 4 ++-- 3 files changed, 8 insertions(+), 8 deletions(-) diff --git a/tests/python/contrib/test_hexagon/test_meta_schedule.py b/tests/python/contrib/test_hexagon/test_meta_schedule.py index 5105bf2749c8..0e772d382d2b 100644 --- a/tests/python/contrib/test_hexagon/test_meta_schedule.py +++ b/tests/python/contrib/test_hexagon/test_meta_schedule.py @@ -64,7 +64,7 @@ def main( # type: ignore # pylint: disable=no-self-argument @tvm.testing.requires_hexagon def test_builder_runner(hexagon_launcher): - if android_serial_number() == "simulator": + if android_serial_number()[0] == "simulator": pytest.skip(msg="Tuning on simulator not supported.") mod = MatmulModule @@ -177,7 +177,7 @@ def verify_dense(sch, target, M, N, K, hexagon_session): @tvm.testing.requires_hexagon def test_vrmpy_dense(hexagon_launcher): - if android_serial_number() == "simulator": + if android_serial_number()[0] == "simulator": pytest.skip(msg="Tuning on simulator not supported.") do_tune = True @@ -275,7 +275,7 @@ def main( # type: ignore @tvm.testing.requires_hexagon def test_vrmpy_dense_auto_tensorize(hexagon_launcher): - if android_serial_number() == "simulator": + if android_serial_number()[0] == "simulator": pytest.skip(msg="Tuning on simulator not supported.") M, N, K = 128, 768, 768 @@ -338,7 +338,7 @@ def test_vrmpy_dense_auto_tensorize(hexagon_launcher): @tvm.testing.requires_hexagon def test_conv2d_relay_auto_schedule(hexagon_launcher): - if android_serial_number() == "simulator": + if android_serial_number()[0] == "simulator": pytest.skip(msg="Tuning on simulator not supported.") I, O, H, W = 64, 64, 56, 56 @@ -418,7 +418,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 android_serial_number() == "simulator": + if android_serial_number()[0] == "simulator": pytest.skip(msg="Tuning on simulator not supported.") target_hexagon = tvm.target.hexagon("v69") diff --git a/tests/python/contrib/test_hexagon/test_models.py b/tests/python/contrib/test_hexagon/test_models.py index 7d60072a0367..d1554eb70e91 100644 --- a/tests/python/contrib/test_hexagon/test_models.py +++ b/tests/python/contrib/test_hexagon/test_models.py @@ -91,7 +91,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 android_serial_number() == "simulator": + if android_serial_number()[0] == "simulator": pytest.skip(msg="Skip on simulator due to long runtime.") dtype = "float32" 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 5e53da92c49e..5f0f0e6b8a95 100644 --- a/tests/python/contrib/test_hexagon/topi/test_cast_slice.py +++ b/tests/python/contrib/test_hexagon/topi/test_cast_slice.py @@ -76,7 +76,7 @@ def test_cast_fp16_fp32_slice( """ Top level testing function for cast fp16 to fp32 """ - if android_serial_number() != "simulator": + if android_serial_number()[0] != "simulator": pytest.skip(msg="Due to https://github.com/apache/tvm/issues/11957") cast_input = te.placeholder(input_shape, name="A", dtype=dtype) @@ -162,7 +162,7 @@ def test_cast_fp32_fp16_slice( """ Top level testing function for cast fp32 to fp16 """ - if android_serial_number() != "simulator": + if android_serial_number()[0] != "simulator": pytest.skip(msg="Due to https://github.com/apache/tvm/issues/11957") cast_input = te.placeholder(input_shape, name="A", dtype=dtype) From a1020d375f6d265a91df045b83095340b732ac1c Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Fri, 21 Oct 2022 06:09:44 +0000 Subject: [PATCH 05/13] fix test on device --- .../python/contrib/test_hexagon/test_benchmark_elemwise_add.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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 769e9bbfce87..ee58514569ae 100644 --- a/tests/python/contrib/test_hexagon/test_benchmark_elemwise_add.py +++ b/tests/python/contrib/test_hexagon/test_benchmark_elemwise_add.py @@ -244,7 +244,7 @@ def _benchmark_hexagon_elementwise_add_kernel( # 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, hexagon_session + target_dso_binary_pathname ) # Create the target-side tensors to hold the primfunc's inputs and outputs... From 482ee94c505a8db2869f468a928cbe82823f015a Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Fri, 21 Oct 2022 18:13:26 +0000 Subject: [PATCH 06/13] add serial device to rpc key --- python/tvm/contrib/hexagon/build.py | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/python/tvm/contrib/hexagon/build.py b/python/tvm/contrib/hexagon/build.py index 7582dcb7becc..1be76139c64c 100644 --- a/python/tvm/contrib/hexagon/build.py +++ b/python/tvm/contrib/hexagon/build.py @@ -264,6 +264,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 != "" + 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_ = [] @@ -271,7 +273,7 @@ def __init__( self._clear_logcat = clear_logcat self._sysmon_profile = sysmon_profile self._sysmon_process = None - rpc_info["device_key"] = HEXAGON_REMOTE_DEVICE_KEY + rpc_info["device_key"] = HEXAGON_REMOTE_DEVICE_KEY + "." + self._serial_number super(HexagonLauncherAndroid, self).__init__(rpc_info, workspace) @@ -556,13 +558,14 @@ 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" + super(HexagonLauncherSimulator, self).__init__(rpc_info, workspace) + def _copy_to_remote( self, local_path: Union[str, pathlib.Path], remote_path: Union[str, pathlib.Path] ): From 2115a5f9e9b322e7bf59fa5d205a1af8c8f45b62 Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Fri, 21 Oct 2022 18:44:55 +0000 Subject: [PATCH 07/13] update error name --- python/tvm/contrib/hexagon/build.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/tvm/contrib/hexagon/build.py b/python/tvm/contrib/hexagon/build.py index 1be76139c64c..cdba63309078 100644 --- a/python/tvm/contrib/hexagon/build.py +++ b/python/tvm/contrib/hexagon/build.py @@ -264,7 +264,7 @@ 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 != "" + 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] From 371c2d38c65343d710cc8de73ae33a61c115c567 Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Sat, 22 Oct 2022 02:44:51 +0000 Subject: [PATCH 08/13] fix comment --- python/tvm/contrib/hexagon/pytest_plugin.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/tvm/contrib/hexagon/pytest_plugin.py b/python/tvm/contrib/hexagon/pytest_plugin.py index 6b8bb61e8b0d..7cbdbec15643 100644 --- a/python/tvm/contrib/hexagon/pytest_plugin.py +++ b/python/tvm/contrib/hexagon/pytest_plugin.py @@ -220,7 +220,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() From 311611578f6223a840eaa2f298abfa7bba0809af Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Mon, 24 Oct 2022 16:33:53 +0000 Subject: [PATCH 09/13] move create_session to launcher --- python/tvm/contrib/hexagon/build.py | 26 +++++++++++++++++-- python/tvm/contrib/hexagon/meta_schedule.py | 3 +-- python/tvm/contrib/hexagon/pytest_plugin.py | 4 +-- .../python/contrib/test_hexagon/README_RPC.md | 2 +- .../metaschedule_e2e/test_resnet50_fp16.py | 3 +-- .../metaschedule_e2e/test_resnet50_int8.py | 2 +- .../test_hexagon/test_meta_schedule.py | 9 +++---- .../test_software_pipeline_async.py | 3 +-- 8 files changed, 35 insertions(+), 17 deletions(-) diff --git a/python/tvm/contrib/hexagon/build.py b/python/tvm/contrib/hexagon/build.py index cdba63309078..321f36f829d9 100644 --- a/python/tvm/contrib/hexagon/build.py +++ b/python/tvm/contrib/hexagon/build.py @@ -34,7 +34,7 @@ import tvm from tvm.contrib.hexagon.hexagon_profiler import HexagonProfiler from ..._ffi import libinfo - +from .session import Session HEXAGON_RPC_LIB_DIR = os.environ.get("HEXAGON_RPC_LIB_DIR") ANDROID_BASH_FILE_NAME = "android_bash.sh" @@ -103,7 +103,7 @@ class HexagonLauncherRPC(metaclass=abc.ABCMeta): The basic flow of interaction with the launcher is launcher = HexagonLauncher(...) launcher.start_server() - with session.create_session() as session: + with launcher.create_session() as session: # Do something with the session launcher.stop_server() @@ -221,6 +221,28 @@ def get_profile_output( Path of the profiling data file """ ... + 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_remote_kw = { + "host": self._rpc_info["rpc_tracker_host"], + "port": self._rpc_info["rpc_tracker_port"], + "priority": 0, + "timeout": 0, + "key": self._rpc_info["device_key"], + } + return Session(self._workspace, hexagon_remote_kw, session_name=session_name) class HexagonLauncherAndroid(HexagonLauncherRPC): diff --git a/python/tvm/contrib/hexagon/meta_schedule.py b/python/tvm/contrib/hexagon/meta_schedule.py index 7ea2c9c272a7..aaf3f8c7f8d5 100644 --- a/python/tvm/contrib/hexagon/meta_schedule.py +++ b/python/tvm/contrib/hexagon/meta_schedule.py @@ -35,7 +35,6 @@ from .build import HexagonLauncherRPC from .tools import export_module -from .session import create_session @derived_object @@ -101,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 create_session(hexagon_launcher._workspace, hexagon_launcher._rpc_info) 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 7cbdbec15643..d16774ea3eaa 100644 --- a/python/tvm/contrib/hexagon/pytest_plugin.py +++ b/python/tvm/contrib/hexagon/pytest_plugin.py @@ -28,7 +28,7 @@ import tvm import tvm.rpc.tracker from tvm.contrib.hexagon.build import HexagonLauncher, HexagonLauncherRPC -from tvm.contrib.hexagon.session import Session, create_session +from tvm.contrib.hexagon.session import Session HEXAGON_TOOLCHAIN = "HEXAGON_TOOLCHAIN" TVM_TRACKER_HOST = "TVM_TRACKER_HOST" @@ -274,7 +274,7 @@ def hexagon_session(hexagon_launcher: HexagonLauncherRPC) -> Session: if hexagon_launcher is None: yield None else: - with create_session(hexagon_launcher._workspace, hexagon_launcher._rpc_info) as session: + with hexagon_launcher.create_session() as session: yield session diff --git a/tests/python/contrib/test_hexagon/README_RPC.md b/tests/python/contrib/test_hexagon/README_RPC.md index 4aff7c203e70..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 `session.create_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/metaschedule_e2e/test_resnet50_fp16.py b/tests/python/contrib/test_hexagon/metaschedule_e2e/test_resnet50_fp16.py index 2f16d6050bf4..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,7 +25,6 @@ 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 tvm.contrib.hexagon.session import create_session from ..infrastructure import get_hexagon_target @@ -105,7 +104,7 @@ def test_resnet50(hexagon_launcher): llvm_graph_mod.run() ref_result = llvm_graph_mod.get_output(0).numpy() - with create_session(hexagon_launcher._workspace, hexagon_launcher._rpc_info) 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 8fac962b9b53..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 create_session(hexagon_launcher._workspace, hexagon_launcher._rpc_info) 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_meta_schedule.py b/tests/python/contrib/test_hexagon/test_meta_schedule.py index 0e772d382d2b..ecc8ab7792ee 100644 --- a/tests/python/contrib/test_hexagon/test_meta_schedule.py +++ b/tests/python/contrib/test_hexagon/test_meta_schedule.py @@ -36,7 +36,6 @@ from tvm.tir import FloatImm from tvm.tir.tensor_intrin.hexagon import VRMPY_u8u8i32_INTRIN from tvm.contrib.hexagon.pytest_plugin import android_serial_number -from tvm.contrib.hexagon.session import create_session from .infrastructure import get_hexagon_target @@ -215,7 +214,7 @@ def schedule_dense_for_tune(sch): ) sch = ms.tir_integration.compile_tir(database, workload, target) - with create_session(hexagon_launcher._workspace, hexagon_launcher._rpc_info) as session: + with hexagon_launcher.create_session() as session: verify_dense(sch, get_hexagon_target("v68"), M, N, K, session) @@ -332,7 +331,7 @@ def test_vrmpy_dense_auto_tensorize(hexagon_launcher): else: sch = tvm.tir.Schedule(Module_vrmpy_auto_tensorize, debug_mask="all") - with create_session(hexagon_launcher._workspace, hexagon_launcher._rpc_info) as session: + with hexagon_launcher.create_session() as session: verify_dense(sch, get_hexagon_target("v68"), M, N, K, session) @@ -399,7 +398,7 @@ def test_conv2d_relay_auto_schedule(hexagon_launcher): target=target, ) - with create_session(hexagon_launcher._workspace, hexagon_launcher._rpc_info) as session: + with hexagon_launcher.create_session() as session: rt_mod = session.get_executor_from_factory(lib) rt_mod.set_input("data", data_np) @@ -458,7 +457,7 @@ def test_dense_relay_auto_schedule(hexagon_launcher): target=target, ) - with create_session(hexagon_launcher._workspace, hexagon_launcher._rpc_info) 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_software_pipeline_async.py b/tests/python/contrib/test_hexagon/test_software_pipeline_async.py index 348a3c799eb0..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 create_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 create_session(hexagon_launcher._workspace, hexagon_launcher._rpc_info) 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) From 3700a8687403eb5b9976dfdb80f5a904faaa982a Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Mon, 24 Oct 2022 16:59:19 +0000 Subject: [PATCH 10/13] add is_simulator --- python/tvm/contrib/hexagon/build.py | 10 ++++++++-- python/tvm/contrib/hexagon/pytest_plugin.py | 11 ++++++----- python/tvm/contrib/hexagon/session.py | 5 ++++- python/tvm/contrib/hexagon/tools.py | 2 ++ tests/python/contrib/test_hexagon/benchmark_util.py | 4 +++- .../python/contrib/test_hexagon/test_meta_schedule.py | 10 +++++----- tests/python/contrib/test_hexagon/test_models.py | 2 +- .../contrib/test_hexagon/topi/test_cast_slice.py | 4 ++-- 8 files changed, 31 insertions(+), 17 deletions(-) diff --git a/python/tvm/contrib/hexagon/build.py b/python/tvm/contrib/hexagon/build.py index 321f36f829d9..9f5f55587052 100644 --- a/python/tvm/contrib/hexagon/build.py +++ b/python/tvm/contrib/hexagon/build.py @@ -35,6 +35,7 @@ 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" @@ -133,6 +134,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._serial_number @abc.abstractmethod def start_server(self): @@ -241,9 +243,13 @@ def create_session(self, session_name: str = "hexagon-rpc") -> Session: "priority": 0, "timeout": 0, "key": self._rpc_info["device_key"], + "serial_number": self._serial_number, } return Session(self._workspace, hexagon_remote_kw, session_name=session_name) + def is_simulator(self): + return self._serial_number == HEXAGON_SIMULATOR_NAME + class HexagonLauncherAndroid(HexagonLauncherRPC): """Hexagon Launcher for Android.""" @@ -584,7 +590,7 @@ def __init__(self, rpc_info: dict, workspace: Union[str, pathlib.Path] = None): 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) @@ -714,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/pytest_plugin.py b/python/tvm/contrib/hexagon/pytest_plugin.py index d16774ea3eaa..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 @@ -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() @@ -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 47dee1d29606..a8c703d82d82 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: @@ -121,6 +121,9 @@ def device(self): return self._device + def is_simulator(self): + return self._remote_kw["serial_number"] == HEXAGON_SIMULATOR_NAME + def get_function(self, name): return self._rpc.get_function(name) 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/tests/python/contrib/test_hexagon/benchmark_util.py b/tests/python/contrib/test_hexagon/benchmark_util.py index 8ffdfe5090ca..0ccbe514326c 100644 --- a/tests/python/contrib/test_hexagon/benchmark_util.py +++ b/tests/python/contrib/test_hexagon/benchmark_util.py @@ -23,6 +23,8 @@ import pytest +from tvm.contrib.hexagon.tools import HEXAGON_SIMULATOR_NAME + def skip_benchmarks_flag_and_reason(): """ @@ -37,7 +39,7 @@ def skip_benchmarks_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/test_meta_schedule.py b/tests/python/contrib/test_hexagon/test_meta_schedule.py index ecc8ab7792ee..d474aa79d8cd 100644 --- a/tests/python/contrib/test_hexagon/test_meta_schedule.py +++ b/tests/python/contrib/test_hexagon/test_meta_schedule.py @@ -63,7 +63,7 @@ def main( # type: ignore # pylint: disable=no-self-argument @tvm.testing.requires_hexagon def test_builder_runner(hexagon_launcher): - if android_serial_number()[0] == "simulator": + if hexagon_launcher.is_simulator(): pytest.skip(msg="Tuning on simulator not supported.") mod = MatmulModule @@ -176,7 +176,7 @@ def verify_dense(sch, target, M, N, K, hexagon_session): @tvm.testing.requires_hexagon def test_vrmpy_dense(hexagon_launcher): - if android_serial_number()[0] == "simulator": + if hexagon_launcher.is_simulator(): pytest.skip(msg="Tuning on simulator not supported.") do_tune = True @@ -274,7 +274,7 @@ def main( # type: ignore @tvm.testing.requires_hexagon def test_vrmpy_dense_auto_tensorize(hexagon_launcher): - if android_serial_number()[0] == "simulator": + if hexagon_launcher.is_simulator(): pytest.skip(msg="Tuning on simulator not supported.") M, N, K = 128, 768, 768 @@ -337,7 +337,7 @@ def test_vrmpy_dense_auto_tensorize(hexagon_launcher): @tvm.testing.requires_hexagon def test_conv2d_relay_auto_schedule(hexagon_launcher): - if android_serial_number()[0] == "simulator": + if hexagon_launcher.is_simulator(): pytest.skip(msg="Tuning on simulator not supported.") I, O, H, W = 64, 64, 56, 56 @@ -417,7 +417,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 android_serial_number()[0] == "simulator": + if hexagon_launcher.is_simulator(): pytest.skip(msg="Tuning on simulator not supported.") target_hexagon = tvm.target.hexagon("v69") diff --git a/tests/python/contrib/test_hexagon/test_models.py b/tests/python/contrib/test_hexagon/test_models.py index d1554eb70e91..fd603b2b90d6 100644 --- a/tests/python/contrib/test_hexagon/test_models.py +++ b/tests/python/contrib/test_hexagon/test_models.py @@ -91,7 +91,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 android_serial_number()[0] == "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/topi/test_cast_slice.py b/tests/python/contrib/test_hexagon/topi/test_cast_slice.py index 5f0f0e6b8a95..abe4d92d8f66 100644 --- a/tests/python/contrib/test_hexagon/topi/test_cast_slice.py +++ b/tests/python/contrib/test_hexagon/topi/test_cast_slice.py @@ -76,7 +76,7 @@ def test_cast_fp16_fp32_slice( """ Top level testing function for cast fp16 to fp32 """ - if android_serial_number()[0] != "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) @@ -162,7 +162,7 @@ def test_cast_fp32_fp16_slice( """ Top level testing function for cast fp32 to fp16 """ - if android_serial_number()[0] != "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) From 841e9ff2825680a1df96b5aafceced9853160362 Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Mon, 24 Oct 2022 17:35:53 +0000 Subject: [PATCH 11/13] lint --- python/tvm/contrib/hexagon/build.py | 10 ++++++---- .../python/contrib/test_hexagon/test_meta_schedule.py | 1 - tests/python/contrib/test_hexagon/test_models.py | 1 - .../contrib/test_hexagon/topi/test_cast_slice.py | 1 - 4 files changed, 6 insertions(+), 7 deletions(-) diff --git a/python/tvm/contrib/hexagon/build.py b/python/tvm/contrib/hexagon/build.py index 9f5f55587052..adad7b1a24b0 100644 --- a/python/tvm/contrib/hexagon/build.py +++ b/python/tvm/contrib/hexagon/build.py @@ -125,7 +125,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, @@ -134,7 +136,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._serial_number + self._serial_number = serial_number @abc.abstractmethod def start_server(self): @@ -303,7 +305,7 @@ def __init__( 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] @@ -592,7 +594,7 @@ def __init__(self, rpc_info: dict, workspace: Union[str, pathlib.Path] = None): raise RuntimeError("Please set HEXAGON_TOOLCHAIN env variable") self._serial_number = HEXAGON_SIMULATOR_NAME - super(HexagonLauncherSimulator, self).__init__(rpc_info, workspace) + 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] diff --git a/tests/python/contrib/test_hexagon/test_meta_schedule.py b/tests/python/contrib/test_hexagon/test_meta_schedule.py index d474aa79d8cd..a7f4cbc39cb1 100644 --- a/tests/python/contrib/test_hexagon/test_meta_schedule.py +++ b/tests/python/contrib/test_hexagon/test_meta_schedule.py @@ -35,7 +35,6 @@ from tvm.script import tir as T from tvm.tir import FloatImm from tvm.tir.tensor_intrin.hexagon import VRMPY_u8u8i32_INTRIN -from tvm.contrib.hexagon.pytest_plugin import android_serial_number from .infrastructure import get_hexagon_target diff --git a/tests/python/contrib/test_hexagon/test_models.py b/tests/python/contrib/test_hexagon/test_models.py index fd603b2b90d6..db578d1057a6 100644 --- a/tests/python/contrib/test_hexagon/test_models.py +++ b/tests/python/contrib/test_hexagon/test_models.py @@ -24,7 +24,6 @@ from tvm import relay from tvm.contrib.hexagon.session import Session from tvm.relay.backend import Executor, Runtime -from tvm.contrib.hexagon.pytest_plugin import android_serial_number from .infrastructure import get_hexagon_target 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 abe4d92d8f66..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,7 +22,6 @@ import tvm.testing from tvm import te import tvm.topi.hexagon.slice_ops as sl -from tvm.contrib.hexagon.pytest_plugin import android_serial_number from ..infrastructure import allocate_hexagon_array, transform_numpy, get_hexagon_target From 0741339dddd7dc07dad9d72ff97df4ca17ba6ff2 Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Mon, 24 Oct 2022 22:04:18 +0000 Subject: [PATCH 12/13] address Eric comment on Session object --- python/tvm/contrib/hexagon/build.py | 14 +++++------ python/tvm/contrib/hexagon/session.py | 34 ++++++++++++++++++++------- 2 files changed, 31 insertions(+), 17 deletions(-) diff --git a/python/tvm/contrib/hexagon/build.py b/python/tvm/contrib/hexagon/build.py index adad7b1a24b0..a358de834e82 100644 --- a/python/tvm/contrib/hexagon/build.py +++ b/python/tvm/contrib/hexagon/build.py @@ -238,16 +238,14 @@ def create_session(self, session_name: str = "hexagon-rpc") -> Session: 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._rpc_info["device_key"], + 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(self._workspace, hexagon_remote_kw, session_name=session_name) + return Session(**hexagon_session_kw) def is_simulator(self): return self._serial_number == HEXAGON_SIMULATOR_NAME diff --git a/python/tvm/contrib/hexagon/session.py b/python/tvm/contrib/hexagon/session.py index a8c703d82d82..466103f6e2c9 100644 --- a/python/tvm/contrib/hexagon/session.py +++ b/python/tvm/contrib/hexagon/session.py @@ -38,8 +38,17 @@ class Session: Parameters ---------- - remote_kw : dict - Remote configs for RPC tracker. + remote_workspace : Union[str, pathlib.Path] + Remote workspace path + + 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. @@ -47,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, remote_workspace: Union[str, pathlib.Path], - remote_kw: dict, + 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._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 @@ -71,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, @@ -122,7 +138,7 @@ def device(self): return self._device def is_simulator(self): - return self._remote_kw["serial_number"] == HEXAGON_SIMULATOR_NAME + return self._serial_number == HEXAGON_SIMULATOR_NAME def get_function(self, name): return self._rpc.get_function(name) From 7f290698006d10e4fe9bd631fe144ae978e77b61 Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Tue, 25 Oct 2022 19:59:19 +0000 Subject: [PATCH 13/13] rebase with main --- python/tvm/contrib/hexagon/build.py | 2 +- src/runtime/hexagon/rpc/hexagon/rpc_server.cc | 2 ++ src/runtime/hexagon/rpc/simulator/rpc_server.cc | 2 ++ 3 files changed, 5 insertions(+), 1 deletion(-) diff --git a/python/tvm/contrib/hexagon/build.py b/python/tvm/contrib/hexagon/build.py index a358de834e82..bc3b065dd941 100644 --- a/python/tvm/contrib/hexagon/build.py +++ b/python/tvm/contrib/hexagon/build.py @@ -31,7 +31,6 @@ import tempfile from typing import Union -import tvm from tvm.contrib.hexagon.hexagon_profiler import HexagonProfiler from ..._ffi import libinfo from .session import Session @@ -225,6 +224,7 @@ def get_profile_output( Path of the profiling data file """ ... + def create_session(self, session_name: str = "hexagon-rpc") -> Session: """Create an RPC session. diff --git a/src/runtime/hexagon/rpc/hexagon/rpc_server.cc b/src/runtime/hexagon/rpc/hexagon/rpc_server.cc index 77c3b8e94d40..f39944615bfd 100644 --- a/src/runtime/hexagon/rpc/hexagon/rpc_server.cc +++ b/src/runtime/hexagon/rpc/hexagon/rpc_server.cc @@ -342,6 +342,8 @@ TVM_REGISTER_GLOBAL("tvm.hexagon.get_profile_output") HEXAGON_PRINT(ERROR, "ERROR: Unsupported profiling mode: %s", profiling_mode.c_str()); *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; diff --git a/src/runtime/hexagon/rpc/simulator/rpc_server.cc b/src/runtime/hexagon/rpc/simulator/rpc_server.cc index 2dd36ee287db..f4370bd3c88c 100644 --- a/src/runtime/hexagon/rpc/simulator/rpc_server.cc +++ b/src/runtime/hexagon/rpc/simulator/rpc_server.cc @@ -349,6 +349,8 @@ TVM_REGISTER_GLOBAL("tvm.hexagon.get_profile_output") HEXAGON_PRINT(ERROR, "ERROR: Unsupported profiling mode: %s", profiling_mode.c_str()); *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;