From 6f2eec9f83009d67624b82838d06dfe7e50ad16f Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Wed, 1 Dec 2021 09:08:44 -0800 Subject: [PATCH 01/15] Add Hexagon RPC --- CMakeLists.txt | 1 + apps/cpp_rpc/CMakeLists.txt | 12 +- cmake/libs/hexagon_rpc_skel/CMakeLists.txt | 119 ++++++++ cmake/modules/Hexagon.cmake | 143 ++++++++- python/tvm/contrib/hexagon/__init__.py | 16 + python/tvm/contrib/hexagon/build.py | 279 ++++++++++++++++++ python/tvm/contrib/{ => hexagon}/hexagon.py | 12 +- python/tvm/contrib/hexagon/session.py | 71 +++++ src/runtime/hexagon/rpc/android/session.cc | 119 ++++++++ .../hexagon/rpc/android_bash.sh.template | 29 ++ src/runtime/hexagon/rpc/hexagon/rpc_server.cc | 212 +++++++++++++ src/runtime/hexagon/rpc/hexagon_rpc.idl | 28 ++ src/target/llvm/codegen_hexagon.cc | 6 +- src/target/llvm/llvm_common.h | 5 - tests/lint/check_file_type.py | 2 + .../test_hexagon/proxy_rpc/test_matmul.py | 2 +- .../contrib/test_hexagon/rpc/conftest.py | 82 +++++ .../contrib/test_hexagon/rpc/test_launcher.md | 98 ++++++ .../contrib/test_hexagon/rpc/test_launcher.py | 260 ++++++++++++++++ .../unittest/test_target_codegen_hexagon.py | 4 +- 20 files changed, 1475 insertions(+), 25 deletions(-) create mode 100644 cmake/libs/hexagon_rpc_skel/CMakeLists.txt create mode 100644 python/tvm/contrib/hexagon/__init__.py create mode 100644 python/tvm/contrib/hexagon/build.py rename python/tvm/contrib/{ => hexagon}/hexagon.py (95%) create mode 100644 python/tvm/contrib/hexagon/session.py create mode 100644 src/runtime/hexagon/rpc/android/session.cc create mode 100644 src/runtime/hexagon/rpc/android_bash.sh.template create mode 100644 src/runtime/hexagon/rpc/hexagon/rpc_server.cc create mode 100644 src/runtime/hexagon/rpc/hexagon_rpc.idl create mode 100644 tests/python/contrib/test_hexagon/rpc/conftest.py create mode 100644 tests/python/contrib/test_hexagon/rpc/test_launcher.md create mode 100644 tests/python/contrib/test_hexagon/rpc/test_launcher.py diff --git a/CMakeLists.txt b/CMakeLists.txt index abf9f4a2a945..2a4dfbafd244 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -30,6 +30,7 @@ tvm_option(USE_ROCM "Build with ROCM" OFF) tvm_option(ROCM_PATH "The path to rocm" /opt/rocm) tvm_option(USE_HEXAGON_DEVICE "Build with Hexagon device support in TVM runtime" OFF) tvm_option(USE_HEXAGON_SDK "Path to the Hexagon SDK root (required for Hexagon support in TVM runtime or for building TVM runtime for Hexagon)" /path/to/sdk) +tvm_option(USE_HEXAGON_RPC "Enable Hexagon RPC using minRPC implementation over Android." OFF) tvm_option(USE_HEXAGON_LAUNCHER "Build the Hexagon graph launcher application" OFF) tvm_option(USE_HEXAGON_PROXY_RPC "Build the Hexagon Proxy RPC server application" OFF) tvm_option(USE_RPC "Build with RPC" ON) diff --git a/apps/cpp_rpc/CMakeLists.txt b/apps/cpp_rpc/CMakeLists.txt index ccac53fc3ca0..966448929a0a 100644 --- a/apps/cpp_rpc/CMakeLists.txt +++ b/apps/cpp_rpc/CMakeLists.txt @@ -6,6 +6,8 @@ set(TVM_RPC_SOURCES rpc_server.cc ) +set(TVM_RPC_LINKER_LIBS "") + if(WIN32) list(APPEND TVM_RPC_SOURCES win32_process.cc) endif() @@ -43,4 +45,12 @@ target_include_directories( PUBLIC DMLC_PATH ) -target_link_libraries(tvm_rpc tvm_runtime) +if (BUILD_FOR_ANDROID AND USE_HEXAGON_SDK) + find_hexagon_sdk_root("${USE_HEXAGON_SDK}" "${USE_HEXAGON_ARCH}") + link_directories(${HEXAGON_REMOTE_ROOT}) + list(APPEND TVM_RPC_LINKER_LIBS cdsprpc log) +endif() + +list(APPEND TVM_RPC_LINKER_LIBS tvm_runtime) + +target_link_libraries(tvm_rpc ${TVM_RPC_LINKER_LIBS}) diff --git a/cmake/libs/hexagon_rpc_skel/CMakeLists.txt b/cmake/libs/hexagon_rpc_skel/CMakeLists.txt new file mode 100644 index 000000000000..a4756aa4e088 --- /dev/null +++ b/cmake/libs/hexagon_rpc_skel/CMakeLists.txt @@ -0,0 +1,119 @@ +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you under the Apache License, Version 2.0 (the +# "License"); you may not use this file except in compliance +# with the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, +# software distributed under the License is distributed on an +# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +# KIND, either express or implied. See the License for the +# specific language governing permissions and limitations +# under the License. + +cmake_minimum_required(VERSION 3.2) +include(ExternalProject) +project(HexagonRPCSkel C CXX) + +set(TVM_SOURCE_DIR "${CMAKE_CURRENT_SOURCE_DIR}/../../..") +set(TVM_SRC_DIR "${TVM_SOURCE_DIR}/src") + + +include("${TVM_SOURCE_DIR}/cmake/modules/HexagonSDK.cmake") + +find_hexagon_sdk_root("${USE_HEXAGON_SDK}" "${USE_HEXAGON_ARCH}") + +include_directories(SYSTEM ${HEXAGON_SDK_INCLUDES} ${HEXAGON_REMOTE_ROOT}) + +set(HEXAGON_RPC_H "hexagon_rpc.h") +set(HEXAGON_RPC_SKEL_C "hexagon_rpc_skel.c") +set(HEXAGON_RPC_STUB_C "hexagon_rpc_stub.c") + +include_directories( + "${TVM_SOURCE_DIR}/include" + "${TVM_SOURCE_DIR}/3rdparty/dlpack/include" + "${TVM_SOURCE_DIR}/3rdparty/dmlc-core/include" +) + +set(QAIC_EXE "${HEXAGON_QAIC_EXE}") +foreach(INCDIR IN LISTS HEXAGON_SDK_INCLUDES HEXAGON_REMOTE_ROOT) + list(APPEND QAIC_FLAGS "-I${INCDIR}") +endforeach() + +add_custom_command( + OUTPUT ${HEXAGON_RPC_SKEL_C} ${HEXAGON_RPC_H} + COMMAND ${QAIC_EXE} ${QAIC_FLAGS} "${TVM_SRC_DIR}/runtime/hexagon/rpc/hexagon_rpc.idl" + MAIN_DEPENDENCY "${TVM_SRC_DIR}/runtime/hexagon/rpc/hexagon_rpc.idl" +) + +include_directories(SYSTEM + ${HEXAGON_QURT_INCLUDES} + ${CMAKE_CURRENT_BINARY_DIR} # Output of qaic will go here +) + +link_directories(${HEXAGON_QURT_LIBS}) + +add_definitions(-D_MACH_I32=int) +add_definitions(-DDMLC_CXX11_THREAD_LOCAL=0) +add_definitions(-DDMLC_USE_LOGGING_LIBRARY=) + +# Extra compile flags (both C and C++). +set(EXTRA_COMP_FLAGS + "-O3" + "-m${USE_HEXAGON_ARCH}" +) +string(REGEX REPLACE ";" " " EXTRA_COMP_FLAGS_STR "${EXTRA_COMP_FLAGS}") +set(CMAKE_C_FLAGS "${EXTRA_COMP_FLAGS_STR} ${CMAKE_C_FLAGS}") +set(CMAKE_CXX_FLAGS "${EXTRA_COMP_FLAGS_STR} ${CMAKE_CXX_FLAGS}") + +set(SKEL_SRCS + "${TVM_SRC_DIR}/runtime/hexagon/rpc/hexagon/rpc_server.cc" +) + +set(MINRPC_SRCS + "${TVM_SRC_DIR}/runtime/minrpc/minrpc_server.h" + "${TVM_SRC_DIR}/runtime/minrpc/rpc_reference.h" +) + +set(TVM_RPC_SRC + "${TVM_SRC_DIR}/runtime/rpc/rpc_module.cc" + "${TVM_SRC_DIR}/runtime/rpc/rpc_endpoint.cc" + "${TVM_SRC_DIR}/runtime/rpc/rpc_session.cc" + "${TVM_SRC_DIR}/runtime/rpc/rpc_local_session.cc" +) + +add_library(hexagon_rpc_skel SHARED + "${HEXAGON_RPC_H}" + "${HEXAGON_RPC_SKEL_C}" + "${SKEL_SRCS}" + "${MINRPC_SRCS}" + "${TVM_RPC_SRC}" +) + +ExternalProject_Add(static_hexagon_tvm_runtime + SOURCE_DIR "${TVM_SOURCE_DIR}" + BUILD_COMMAND $(MAKE) runtime + CMAKE_ARGS + "-DCMAKE_C_COMPILER=${CMAKE_C_COMPILER}" + "-DCMAKE_CXX_COMPILER=${CMAKE_CXX_COMPILER}" + "-DUSE_HEXAGON_ARCH=${USE_HEXAGON_ARCH}" + "-DCMAKE_CXX_STANDARD=14" + "-DUSE_LIBBACKTRACE=OFF" + "-DUSE_LLVM=OFF" + "-DUSE_RPC=OFF" + "-DBUILD_STATIC_RUNTIME=ON" + "-DUSE_HEXAGON_SDK=${USE_HEXAGON_SDK}" + INSTALL_COMMAND "" + BUILD_ALWAYS ON +) +ExternalProject_Get_Property(static_hexagon_tvm_runtime BINARY_DIR) + +add_dependencies(hexagon_rpc_skel static_hexagon_tvm_runtime) +add_library(h_tvm_runtime STATIC IMPORTED) +set_target_properties(h_tvm_runtime PROPERTIES IMPORTED_LOCATION "${BINARY_DIR}/libtvm_runtime.a") + +target_link_libraries(hexagon_rpc_skel -Wl,--whole-archive h_tvm_runtime -Wl,--no-whole-archive) diff --git a/cmake/modules/Hexagon.cmake b/cmake/modules/Hexagon.cmake index 6fc64cc954dd..6690c1a922ed 100644 --- a/cmake/modules/Hexagon.cmake +++ b/cmake/modules/Hexagon.cmake @@ -53,7 +53,6 @@ if(BUILD_FOR_HEXAGON) include_directories(SYSTEM ${HEXAGON_SDK_INCLUDES} ${HEXAGON_QURT_INCLUDES}) endif() - if (NOT USE_HEXAGON_SDK STREQUAL "" AND NOT USE_HEXAGON_SDK STREQUAL "/path/to/sdk") set(HEXAGON_SDK_PATH_DEFINED ${USE_HEXAGON_SDK}) @@ -73,10 +72,9 @@ endif() # e.g. when compiling the TVM runtime for Hexagon. if (NOT BUILD_FOR_HEXAGON AND NOT BUILD_FOR_ANDROID) if(USE_HEXAGON_LAUNCHER STREQUAL "OFF" AND - USE_HEXAGON_PROXY_RPC STREQUAL "OFF") + USE_HEXAGON_PROXY_RPC STREQUAL "OFF" AND NOT USE_HEXAGON_RPC) if(USE_HEXAGON_DEVICE STREQUAL "OFF") list(APPEND COMPILER_SRCS src/target/opt/build_hexagon_off.cc) - return() elseif(NOT USE_HEXAGON_DEVICE STREQUAL "${PICK_SIM}" AND NOT USE_HEXAGON_DEVICE STREQUAL "${PICK_HW}") set(ERROR_MSG @@ -202,6 +200,108 @@ if(USE_HEXAGON_PROXY_RPC STREQUAL "ON") set_directory_properties(PROPERTIES ADDITIONAL_MAKE_CLEAN_FILES "${RPC_BINARY_DIR}") endif() +if(USE_HEXAGON_RPC) + if(DEFINED USE_ANDROID_TOOLCHAIN) + if(NOT DEFINED ANDROID_PLATFORM) + message(SEND_ERROR "Please set ANDROID_PLATFORM " + "when providing an Android cmake toolchain.") + endif() + if(NOT DEFINED ANDROID_ABI) + message(SEND_ERROR "Please set ANDROID_ABI " + "when providing an Android cmake toolchain.") + endif() + else() + message(SEND_ERROR "Please set USE_ANDROID_TOOLCHAIN to build the android " + "RPC server for Hexagon.") + endif() + + if(NOT DEFINED USE_HEXAGON_SDK) + message(SEND_ERROR "Please set USE_HEXAGON_SDK to build the android " + "RPC server for Hexagon RPC.") + endif() + if(NOT DEFINED USE_HEXAGON_ARCH) + message(SEND_ERROR "Please set USE_HEXAGON_ARCH to build the android " + "RPC server for Hexagon RPC.") + endif() + find_hexagon_sdk_root("${USE_HEXAGON_SDK}" "${USE_HEXAGON_ARCH}") + + set(HEXAGON_RPC_OUTPUT "${CMAKE_CURRENT_BINARY_DIR}/hexagon_rpc") + file(MAKE_DIRECTORY ${HEXAGON_RPC_OUTPUT}) + + # Android Part + ExternalProject_Add(android_runtime_rpc + SOURCE_DIR "${CMAKE_SOURCE_DIR}" + BUILD_COMMAND $(MAKE) runtime tvm_rpc + CMAKE_ARGS + "-DCMAKE_TOOLCHAIN_FILE=${USE_ANDROID_TOOLCHAIN}" + "-DUSE_ANDROID_TOOLCHAIN=${USE_ANDROID_TOOLCHAIN}" + "-DANDROID_PLATFORM=${ANDROID_PLATFORM}" + "-DANDROID_ABI=${ANDROID_ABI}" + "-DCMAKE_CXX_STANDARD=14" + "-DUSE_LIBBACKTRACE=OFF" + "-DUSE_LLVM=OFF" + "-DUSE_RPC=ON" + "-DUSE_CPP_RPC=ON" + "-DUSE_HEXAGON_SDK=${USE_HEXAGON_SDK}" + "-DUSE_HEXAGON_ARCH=${USE_HEXAGON_ARCH}" + "-DCMAKE_VERBOSE_MAKEFILE=ON" + INSTALL_COMMAND "" + BUILD_ALWAYS ON + ) + ExternalProject_Get_Property(android_runtime_rpc BINARY_DIR) + ExternalProject_Add_Step(android_runtime_rpc copy_binary_runtime + COMMAND ${CMAKE_COMMAND} -E copy_if_different + ${BINARY_DIR}/libtvm_runtime.so + ${HEXAGON_RPC_OUTPUT}/libtvm_runtime_android.so + DEPENDEES install + ) + ExternalProject_Add_Step(android_runtime_rpc copy_binary_rpc + COMMAND ${CMAKE_COMMAND} -E copy_if_different + ${BINARY_DIR}/tvm_rpc + ${HEXAGON_RPC_OUTPUT}/tvm_rpc_android + DEPENDEES install + ) + + if("${USE_HEXAGON_TOOLCHAIN}" STREQUAL "") + message(SEND_ERROR "Please set USE_HEXAGON_TOOLCHAIN to build the hexagon " + "RPC SKEL.") + endif() + find_hexagon_toolchain() + message(STATUS "HEXAGON_TOOLCHAIN: ${HEXAGON_TOOLCHAIN}") + + # Hexagon Part + ExternalProject_Add(hexagon_rpc_skel + SOURCE_DIR "${CMAKE_SOURCE_DIR}/cmake/libs/hexagon_rpc_skel" + INSTALL_DIR "${LAUNCHER_BINARY_DIR}" + CMAKE_ARGS + "-DCMAKE_C_COMPILER=${HEXAGON_TOOLCHAIN}/bin/hexagon-clang" + "-DCMAKE_CXX_COMPILER=${HEXAGON_TOOLCHAIN}/bin/hexagon-clang++" + "-DFASTRPC_LIBS=SKEL" + "-DUSE_HEXAGON_ARCH=${USE_HEXAGON_ARCH}" + "-DUSE_HEXAGON_SDK=${USE_HEXAGON_SDK}" + INSTALL_COMMAND "" + BUILD_ALWAYS ON + ) + ExternalProject_Get_Property(hexagon_rpc_skel BINARY_DIR) + ExternalProject_Add_Step(hexagon_rpc_skel copy_hexagon_skel + COMMAND ${CMAKE_COMMAND} -E copy_if_different + ${BINARY_DIR}/libhexagon_rpc_skel.so + ${HEXAGON_RPC_OUTPUT}/libhexagon_rpc_skel.so + DEPENDEES install + ) + + # copy android_bash template file + configure_file("${CMAKE_SOURCE_DIR}/src/runtime/hexagon/rpc/android_bash.sh.template" + ${HEXAGON_RPC_OUTPUT} COPYONLY) +endif() + +if (USE_HEXAGON_SDK AND BUILD_FOR_ANDROID) + find_hexagon_sdk_root("${USE_HEXAGON_SDK}" "${USE_HEXAGON_ARCH}") + include_directories(SYSTEM ${HEXAGON_SDK_INCLUDES} ${HEXAGON_REMOTE_ROOT}) + link_directories(${HEXAGON_REMOTE_ROOT}) + list(APPEND TVM_RUNTIME_LINKER_LIBS cdsprpc) +endif() + if(USE_HEXAGON_DEVICE STREQUAL "${PICK_SIM}") find_hexagon_toolchain() message(STATUS "Hexagon toolchain: ${HEXAGON_TOOLCHAIN}") @@ -217,16 +317,20 @@ if(USE_HEXAGON_DEVICE STREQUAL "${PICK_SIM}") "-DHEXAGON_ARCH=${USE_HEXAGON_ARCH}" INSTALL_COMMAND "true" ) -elseif(USE_HEXAGON_DEVICE STREQUAL "${PICK_HW}") +elseif((USE_HEXAGON_DEVICE STREQUAL "${PICK_HW}") OR (USE_HEXAGON_RPC AND BUILD_FOR_ANDROID)) find_hexagon_sdk_root("${USE_HEXAGON_SDK}" "${USE_HEXAGON_ARCH}") find_hexagon_toolchain() - file(GLOB RUNTIME_HEXAGON_DEVICE_SRCS src/runtime/hexagon/android/target/*.cc) + + if(USE_HEXAGON_DEVICE STREQUAL "${PICK_HW}") + file(GLOB RUNTIME_HEXAGON_DEVICE_SRCS src/runtime/hexagon/android/target/*.cc) + endif() include_directories(SYSTEM ${HEXAGON_SDK_INCLUDES} ${HEXAGON_RPCMEM_ROOT}/inc ${HEXAGON_REMOTE_ROOT} ) + list(APPEND TVM_RUNTIME_LINKER_LIBS "dl") if(BUILD_FOR_ANDROID) # Hexagon runtime uses __android_log_print, which is in liblog. @@ -240,11 +344,36 @@ if (USE_HEXAGON_DEVICE STREQUAL "${PICK_NONE}") file(GLOB RUNTIME_HEXAGON_SRCS src/runtime/hexagon/hexagon/*.cc) elseif(BUILD_FOR_ANDROID AND HEXAGON_SDK_PATH_DEFINED) list(APPEND RUNTIME_HEXAGON_SRCS src/runtime/hexagon/proxy_rpc/device_api.cc) - else() + elseif(USE_HEXAGON_RPC) file(GLOB RUNTIME_HEXAGON_SRCS src/runtime/hexagon/host/*.cc) endif() else() file(GLOB RUNTIME_HEXAGON_SRCS src/runtime/hexagon/android/*.cc) endif() + +if(USE_HEXAGON_SDK AND BUILD_FOR_ANDROID) + find_hexagon_sdk_root("${USE_HEXAGON_SDK}" "${USE_HEXAGON_ARCH}") + include_directories(SYSTEM ${HEXAGON_SDK_INCLUDES} ${HEXAGON_REMOTE_ROOT}) + + set(QAIC_EXE "${HEXAGON_QAIC_EXE}") + foreach(INCDIR IN LISTS HEXAGON_SDK_INCLUDES HEXAGON_REMOTE_ROOT) + list(APPEND QAIC_FLAGS "-I${INCDIR}") + endforeach() + + set(HEXAGON_RPC_DIR "${CMAKE_SOURCE_DIR}/src/runtime/hexagon/rpc") + set(RPC_IDL "hexagon_rpc.idl") + set(RPC_H "hexagon_rpc.h") + set(RPC_STUB_C "hexagon_rpc_stub.c") + + add_custom_command( + OUTPUT "${HEXAGON_RPC_DIR}/${RPC_STUB_C}" "${HEXAGON_RPC_DIR}/${RPC_H}" + COMMAND ${QAIC_EXE} ${QAIC_FLAGS} "${HEXAGON_RPC_DIR}/${RPC_IDL}" -o ${HEXAGON_RPC_DIR} + MAIN_DEPENDENCY "${HEXAGON_RPC_DIR}/${RPC_IDL}" + ) + file(GLOB HEXAGON_RPC_CPP "${HEXAGON_RPC_DIR}/android/*.cc") + set(HEXAGON_RPC_STUB_C "${HEXAGON_RPC_DIR}/${RPC_STUB_C}") +endif() + list(APPEND RUNTIME_SRCS ${RUNTIME_HEXAGON_SRCS} ${RUNTIME_HEXAGON_SIM_SRCS} - ${RUNTIME_HEXAGON_DEVICE_SRCS} ${RUNTIME_HEXAGON_COMMON_SRCS}) + ${RUNTIME_HEXAGON_DEVICE_SRCS} ${HEXAGON_RPC_CPP} ${HEXAGON_RPC_STUB_C} + ${RUNTIME_HEXAGON_COMMON_SRCS}) diff --git a/python/tvm/contrib/hexagon/__init__.py b/python/tvm/contrib/hexagon/__init__.py new file mode 100644 index 000000000000..13a83393a912 --- /dev/null +++ b/python/tvm/contrib/hexagon/__init__.py @@ -0,0 +1,16 @@ +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you under the Apache License, Version 2.0 (the +# "License"); you may not use this file except in compliance +# with the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, +# software distributed under the License is distributed on an +# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +# KIND, either express or implied. See the License for the +# specific language governing permissions and limitations +# under the License. diff --git a/python/tvm/contrib/hexagon/build.py b/python/tvm/contrib/hexagon/build.py new file mode 100644 index 000000000000..393e60b43594 --- /dev/null +++ b/python/tvm/contrib/hexagon/build.py @@ -0,0 +1,279 @@ +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you under the Apache License, Version 2.0 (the +# "License"); you may not use this file except in compliance +# with the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, +# software distributed under the License is distributed on an +# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +# KIND, either express or implied. See the License for the +# specific language governing permissions and limitations +# under the License. + +"""Defines top-level glue functions for building Hexagon.""" + +import pathlib +import os +import subprocess +from typing import Union +import stat +import datetime + +import tvm +from tvm import rpc as _rpc +from ..._ffi import libinfo +from .session import Session + + +RPC_SERVER_FILES = ["tvm_rpc_android", "libtvm_runtime_android.so", "android_bash.sh"] + +HEXAGON_FILES = ["libhexagon_rpc_skel.so"] + +HEXAGON_RPC_DIR = None + +ANDROID_HEXAGON_TEST_BASE_DIR = pathlib.Path("/data/local/tmp/hexagon_test") + + +def get_hexagon_rpc_dir() -> pathlib.Path: + """Find the Hexagon library. + + Returns + ------- + str : + The path to the Hexagon library + """ + global HEXAGON_RPC_DIR + if HEXAGON_RPC_DIR is None: + for path in libinfo.find_lib_path(): + rpc_dir = os.path.join(os.path.dirname(path), "hexagon_rpc") + if os.path.isdir(rpc_dir): + HEXAGON_RPC_DIR = pathlib.Path(rpc_dir) + break + else: + raise "hexagon_rpc was not found." + return HEXAGON_RPC_DIR + + +class HexagonLauncher: + """Hexagon Launcher""" + + def __init__(self, serial_number: str): + """Configure a new HexagonLauncher + + Parameters + ---------- + serial_number : str + Android device serial number from android 'adb' command. + """ + # Hexagon RPCSession + self.session = None + + self._serial_number = serial_number + self._adb_device_sub_cmd = ["adb", "-s", self._serial_number] + self._android_remote = None + self._mod = None + self._workspace = None + + ANDROID_REMOTE_DEVICE_KEY = "hexagon_dev-android" + HEXAGON_REMOTE_DEVICE_KEY = "hexagon_dev-hexagon" + + def android_run_rpc( + self, + workspace_dir: Union[str, pathlib.Path] = None, + rpc_server_port: int = 7070, + rpc_tracker_host: str = "0.0.0.0", + rpc_tracker_port: int = 9190, + ): + """Upload Android artifacts and run RPC server on Android. + + Parameters + ---------- + workspace_dir : Union[str, pathlib.Path] + Workspace directory used on Android to upload artifacts. + + rpc_server_port : int + Android RPC server port number + + rpc_tracker_host : str + RPC tracker IP on host + + rpc_tracker_port : int + RPC tracker port on host + """ + if not workspace_dir: + self._workspace = str( + ANDROID_HEXAGON_TEST_BASE_DIR + / datetime.datetime.now().strftime("%Y-%m-%dT%H-%M-%S") + ) + else: + self._workspace = workspace_dir + + # Upload RPC server and libraries + subprocess.check_call(self._adb_device_sub_cmd + ["shell", "mkdir", "-p", self._workspace]) + + # create bash script + android_bash_script_path = get_hexagon_rpc_dir() / "android_bash.sh" + with open(get_hexagon_rpc_dir() / "android_bash.sh.template", "r") as src_f: + if os.path.exists(android_bash_script_path): + os.remove(android_bash_script_path) + with open(android_bash_script_path, "w") as dest_f: + for line in src_f.readlines(): + if "" in line: + line = line.replace("", str(rpc_tracker_host)) + if "" in line: + line = line.replace("", str(rpc_tracker_port)) + if f"" in line: + line = line.replace( + "", self.ANDROID_REMOTE_DEVICE_KEY + ) + if "" in line: + line = line.replace( + "", self.HEXAGON_REMOTE_DEVICE_KEY + ) + if "" in line: + line = line.replace("", str(rpc_server_port)) + dest_f.write(line) + + # make shell script executable + android_bash_stat = os.stat(android_bash_script_path) + os.chmod(android_bash_script_path, android_bash_stat.st_mode | stat.S_IEXEC) + + # push files + for item in RPC_SERVER_FILES: + src_path = get_hexagon_rpc_dir() / item + destination = f"{self._workspace}/{item}" + if item == "libtvm_runtime_android.so": + destination = f"{self._workspace}/libtvm_runtime.so" + subprocess.check_call(self._adb_device_sub_cmd + ["push", src_path, destination]) + + # enable root for adb + subprocess.check_call(self._adb_device_sub_cmd + ["root"]) + # Removed pre-defined forward/reverse rules + subprocess.check_call(self._adb_device_sub_cmd + ["forward", "--remove-all"]) + subprocess.check_call(self._adb_device_sub_cmd + ["reverse", "--remove-all"]) + + # Enable port reverse for RPC tracker + subprocess.check_call( + self._adb_device_sub_cmd + + ["reverse", f"tcp:{rpc_tracker_port}", f"tcp:{rpc_tracker_port}"] + ) + # Enable port forward for RPC server. We forward 9 ports after the rpc_server_port. + for i in range(0, 10): + subprocess.check_call( + self._adb_device_sub_cmd + + ["forward", f"tcp:{rpc_server_port+i}", f"tcp:{rpc_server_port+i}"] + ) + + # Run server and connect to tracker + subprocess.Popen( + self._adb_device_sub_cmd + ["shell", f"cd {self._workspace} && ./android_bash.sh"], + stdout=subprocess.PIPE, + stdin=subprocess.PIPE, + stderr=subprocess.PIPE, + ) + + def hexagon_setup(self): + """Upload Hexagon artifacts on Android.""" + for item in HEXAGON_FILES: + src_path = get_hexagon_rpc_dir() / item + dst_path = f"{self._workspace}/{item}" + subprocess.check_call(self._adb_device_sub_cmd + ["push", src_path, dst_path]) + + def android_remote_setup(self, remote_kw: dict): + """Setup remote connection from host to Android RPC server. + + Parameters + ---------- + remote_kw : dict + RPC tracker configs. + """ + tracker = _rpc.connect_tracker(remote_kw["host"], remote_kw["port"]) + self.android_remote = tracker.request( + self.ANDROID_REMOTE_DEVICE_KEY, + priority=remote_kw["priority"], + session_timeout=remote_kw["timeout"], + ) + + def hexagon_session_setup(self, remote_kw: dict): + """Setup Hexagon RPC Session from host to Hexagon device. + + Parameters + ---------- + remote_kw : dict + RPC tracker configs. + """ + assert self.android_remote, "android_remote should initialied first." + hexagon_remote_kw = dict(remote_kw) + hexagon_remote_kw["key"] = self.HEXAGON_REMOTE_DEVICE_KEY + self.session = Session(hexagon_remote_kw) + + def get_module(self, module_name: str): + """Load a Hexagon TVM module, already uploaded on Android, on Hexagon and return the module. + + Parameters + ---------- + module_name : str + Module filename. + + Returns + ------- + TVMModule : + A TVM Module loaded on hexagon. + """ + module_path = f"{self._workspace}/{module_name}" + self._mod = self.session.load_module(module_path) + return self._mod + + def upload(self, host_path: Union[str, pathlib.Path], remote_filename: str = None): + """Upload a file to remote(Android). + + Parameters + ---------- + host_path : Union[str, pathlib.Path] + File path on host. + + remote_filename : str + File name on remote(Android). + Returns + ------- + TVMModule : + A TVM Module loaded on hexagon. + """ + src_path = str(host_path) + if remote_filename: + dst_remote_path = f"{self._workspace}/{remote_filename}" + else: + dst_remote_path = None + self.android_remote.upload(src_path, target=dst_remote_path) + + def get_local_graph_executor(self, libmod, remote_libmod_filename: str): + """Create a GraphModule. + + Parameters + ---------- + libmod : tvm.runtime.Module + The module of the corresponding function + + remote_libmod_filename : str + Module filename on remote. Assumed this file lives under self._workspace path. + + Returns + ------- + graph_module : GraphModule + Runtime graph module that can be used to execute the graph. + """ + self.session.__enter__() + hexagon_mod = self.get_module(remote_libmod_filename) + return tvm.contrib.graph_executor.create( + libmod.get_graph_json(), hexagon_mod, self.session.device + ) + + def close(self): + """Close RPC servers""" + subprocess.Popen(self._adb_device_sub_cmd + ["shell", "kill `cat android_rpc_pid.txt`"]) + subprocess.Popen(self._adb_device_sub_cmd + ["shell", "kill `cat hexagon_rpc_pid.txt`"]) diff --git a/python/tvm/contrib/hexagon.py b/python/tvm/contrib/hexagon/hexagon.py similarity index 95% rename from python/tvm/contrib/hexagon.py rename to python/tvm/contrib/hexagon/hexagon.py index fe256163f73c..35136a31f3a9 100644 --- a/python/tvm/contrib/hexagon.py +++ b/python/tvm/contrib/hexagon/hexagon.py @@ -22,7 +22,7 @@ import tvm import tvm.ir import tvm.contrib.cc as cc -from .._ffi.registry import register_func +from ..._ffi.registry import register_func # Linking Hexagon shared libraries. @@ -47,16 +47,16 @@ def register_linker(f): """Register a function that will return the path to the Hexagon linker.""" - return register_func("tvm.contrib.hexagon.hexagon_link", f, True) + return register_func("tvm.contrib.hexagon.hexagon.hexagon_link", f, True) -@register_func("tvm.contrib.hexagon.hexagon_link") +@register_func("tvm.contrib.hexagon.hexagon.hexagon_link") def hexagon_link(): """Return path to the Hexagon linker.""" return hexagon_link_main -@register_func("tvm.contrib.hexagon.link_shared") +@register_func("tvm.contrib.hexagon.hexagon.link_shared") def link_shared(so_name, objs, **kwargs): """Link shared library on Hexagon using the registered Hexagon linker. @@ -83,9 +83,9 @@ def to_str(s): objs = [to_str(s) for s in objs] - linker = tvm.get_global_func("tvm.contrib.hexagon.hexagon_link")() + linker = tvm.get_global_func("tvm.contrib.hexagon.hexagon.hexagon_link")() if kwargs.get("verbose"): - print("tvm.contrib.hexagon.link_shared:") + print("tvm.contrib.hexagon.hexagon.link_shared:") print(" Using linker:", linker) print(" Library name:", so_name) print(" Object files:", objs) diff --git a/python/tvm/contrib/hexagon/session.py b/python/tvm/contrib/hexagon/session.py new file mode 100644 index 000000000000..8cf449cbcc5a --- /dev/null +++ b/python/tvm/contrib/hexagon/session.py @@ -0,0 +1,71 @@ +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you under the Apache License, Version 2.0 (the +# "License"); you may not use this file except in compliance +# with the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, +# software distributed under the License is distributed on an +# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +# KIND, either express or implied. See the License for the +# specific language governing permissions and limitations +# under the License. + +"""Defines a Session class for Hexagon devices.""" + +from tvm import rpc as _rpc + + +class Session: + """Hexagon Device Session + + Parameters + ---------- + remote_kw : dict + Remote configs for RPC tracker. + + session_name : str + Hexagon RPC session name. + """ + + def __init__( + self, + remote_kw: dict, + session_name: str = "hexagon-rpc", + ): + self._session_name = session_name + self._remote_kw = remote_kw + self._rpc = None + self.device = None + + def __enter__(self): + if self.device: + return self + + tracker = _rpc.connect_tracker(self._remote_kw["host"], self._remote_kw["port"]) + try: + self._rpc = tracker.request( + self._remote_kw["key"], + priority=self._remote_kw["priority"], + session_timeout=self._remote_kw["timeout"], + session_constructor_args=[ + "tvm.contrib.hexagon.create_hexagon_session", + self._session_name, + ], + ) + self.device = self._rpc.hexagon(0) + return self + + except RuntimeError as exception: + raise exception + + def __exit__(self, exc_type, exc_value, exc_traceback): + pass + + def load_module(self, path: str): + assert isinstance(path, str), f"Invalid path type, {type(path)} != str" + return self._rpc.get_function("tvm.hexagon.load_module")(path) diff --git a/src/runtime/hexagon/rpc/android/session.cc b/src/runtime/hexagon/rpc/android/session.cc new file mode 100644 index 000000000000..886c1c7a22da --- /dev/null +++ b/src/runtime/hexagon/rpc/android/session.cc @@ -0,0 +1,119 @@ +/* + * Licensed to the Apache Software Foundation (ASF) under one + * or more contributor license agreements. See the NOTICE file + * distributed with this work for additional information + * regarding copyright ownership. The ASF licenses this file + * to you under the Apache License, Version 2.0 (the + * "License"); you may not use this file except in compliance + * with the License. You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, + * software distributed under the License is distributed on an + * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY + * KIND, either express or implied. See the License for the + * specific language governing permissions and limitations + * under the License. + */ + +/*! + * \file hexagon_session.cc + */ + +#include + +extern "C" { +#include +#include +#include +#include +} + +#include + +#include + +#include "../../../rpc/rpc_channel.h" +#include "../../../rpc/rpc_endpoint.h" +#include "../../../rpc/rpc_session.h" +#include "../hexagon_rpc.h" + +namespace tvm { +namespace runtime { +namespace hexagon { + +class HexagonTransportChannel : public RPCChannel { + public: + explicit HexagonTransportChannel(const std::string& uri) { + if (_handle != AEE_EUNKNOWN) return; + + enable_unsigned_pd(true); + set_remote_stack_size(128 * 1024); + AEEResult rc = hexagon_rpc_open(uri.c_str(), &_handle); + ICHECK(rc == AEE_SUCCESS) << "Hexagon RPC Open failed. URI: " << uri.c_str(); + } + + size_t Send(const void* data, size_t size) override { + ICHECK(_handle != AEE_EUNKNOWN) << "RPC handle is not initialized."; + AEEResult rc = + hexagon_rpc_send(_handle, static_cast(data), static_cast(size)); + ICHECK(rc == AEE_SUCCESS) << "hexagon_rpc_send failed: " << rc; + return size; + } + + size_t Recv(void* data, size_t size) override { + ICHECK(_handle != AEE_EUNKNOWN) << "RPC handle is not initialized."; + int64_t written_size = 0; + AEEResult rc = hexagon_rpc_receive(_handle, static_cast(data), + static_cast(size), &written_size); + ICHECK(rc == AEE_SUCCESS) << "hexagon_rpc_receive failed: " << rc; + return static_cast(written_size); + } + + AEEResult Close() { + if (_handle == AEE_EUNKNOWN) return AEE_SUCCESS; + return hexagon_rpc_close(_handle); + } + + private: + AEEResult set_remote_stack_size(int size) { + remote_rpc_thread_params data; + data.domain = CDSP_DOMAIN_ID; + data.prio = -1; + data.stack_size = size; + AEEResult rc = remote_session_control(FASTRPC_THREAD_PARAMS, &data, sizeof(data)); + if (rc != AEE_SUCCESS) { + LOG(ERROR) << "error setting remote stack size: " << std::hex << rc << '\n'; + } + return rc; + } + + AEEResult enable_unsigned_pd(bool enable) { + remote_rpc_control_unsigned_module data; + data.domain = CDSP_DOMAIN_ID; + data.enable = static_cast(enable); + AEEResult rc = remote_session_control(DSPRPC_CONTROL_UNSIGNED_MODULE, &data, sizeof(data)); + if (rc != AEE_SUCCESS) { + LOG(ERROR) << "Error " << (enable ? "enabling" : "disabling") << " unsigned PD\n"; + } + return rc; + } + + remote_handle64 _handle = AEE_EUNKNOWN; +}; + +TVM_REGISTER_GLOBAL("tvm.contrib.hexagon.create_hexagon_session") + .set_body([](TVMArgs args, TVMRetValue* rv) { + std::string session_name = args[0]; + HexagonTransportChannel* hexagon_channel = + new HexagonTransportChannel(hexagon_rpc_URI CDSP_DOMAIN); + std::unique_ptr channel(hexagon_channel); + auto ep = RPCEndpoint::Create(std::move(channel), session_name, "", NULL); + auto sess = CreateClientSession(ep); + *rv = CreateRPCSessionModule(sess); + }); + +} // namespace hexagon +} // namespace runtime +} // namespace tvm diff --git a/src/runtime/hexagon/rpc/android_bash.sh.template b/src/runtime/hexagon/rpc/android_bash.sh.template new file mode 100644 index 000000000000..f6afa793147c --- /dev/null +++ b/src/runtime/hexagon/rpc/android_bash.sh.template @@ -0,0 +1,29 @@ +#!/bin/sh +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you under the Apache License, Version 2.0 (the +# "License"); you may not use this file except in compliance +# with the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, +# software distributed under the License is distributed on an +# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +# KIND, either express or implied. See the License for the +# specific language governing permissions and limitations +# under the License. + +export LD_LIBRARY_PATH=. +./tvm_rpc_android server --port= --tracker=: --key=& +android_rpc_pid=$! +./tvm_rpc_android server --port= --tracker=: --key=& +hexagon_rpc_pid=$! + +rm -f android_rpc_pid.txt +echo $android_rpc_pid >> android_rpc_pid.txt + +rm -f hexagon_rpc_pid.txt +echo $hexagon_rpc_pid >> hexagon_rpc_pid.txt diff --git a/src/runtime/hexagon/rpc/hexagon/rpc_server.cc b/src/runtime/hexagon/rpc/hexagon/rpc_server.cc new file mode 100644 index 000000000000..b2f04cf4879b --- /dev/null +++ b/src/runtime/hexagon/rpc/hexagon/rpc_server.cc @@ -0,0 +1,212 @@ +/* + * Licensed to the Apache Software Foundation (ASF) under one + * or more contributor license agreements. See the NOTICE file + * distributed with this work for additional information + * regarding copyright ownership. The ASF licenses this file + * to you under the Apache License, Version 2.0 (the + * "License"); you may not use this file except in compliance + * with the License. You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, + * software distributed under the License is distributed on an + * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY + * KIND, either express or implied. See the License for the + * specific language governing permissions and limitations + * under the License. + */ + +extern "C" { +#include +#include +#include +#include +#include +#include +} + +#include +#include +#include +#include + +#include +#include +#include + +#include "../../../library_module.h" +#include "../../../minrpc/minrpc_server.h" +#include "../../hexagon/hexagon_common.h" +#include "hexagon_rpc.h" + +#define TVM_HEXAGON_RPC_BUFF_SIZE_BYTES 1024 * 1024 + +#define TVM_LOG_CUSTOMIZE 1 + +namespace tvm { +namespace runtime { +namespace hexagon { + +class HexagonIOHandler { + public: + explicit HexagonIOHandler(uint8_t* read_buffer) : read_buffer_{read_buffer} {} + + void MessageStart(size_t message_size_bytes) {} + + ssize_t PosixWrite(const uint8_t* buf, size_t write_len_bytes) { + FARF(ALWAYS, "HexagonIOHandler PosixWrite called, write_len_bytes: %d", write_len_bytes); + size_t written_size = static_cast( + write_buffer_.sputn(reinterpret_cast(buf), write_len_bytes)); + if (written_size != write_len_bytes) { + FARF(ALWAYS, "HexagonIOHandler written_size failed"); + } + return (ssize_t)written_size; + } + + void MessageDone() {} + + ssize_t PosixRead(uint8_t* buf, size_t read_len_bytes) { + FARF(ALWAYS, "HexagonIOHandler PosixRead called, %d, %d", read_len_bytes, + read_buffer_size_bytes_); + + uint32_t bytes_to_read = 0; + if ((read_buffer_size_bytes_ - read_len_bytes) < 0) { + bytes_to_read = read_buffer_size_bytes_; + } else { + bytes_to_read = read_len_bytes; + } + + std::memcpy(buf, read_buffer_, bytes_to_read); + read_buffer_ += bytes_to_read; + read_buffer_size_bytes_ -= bytes_to_read; + if (bytes_to_read != read_len_bytes) { + FARF(ERROR, "Error bytes_to_read (%d) < read_len_bytes (%d).", bytes_to_read, read_len_bytes); + } + return (ssize_t)bytes_to_read; + } + + void SetReadBuffer(const uint8_t* buf, size_t buf_size_bytes) { + FARF(ALWAYS, "HexagonIOHandler SetReadBuffer called: %d, prev read_buffer_size_bytes_: ", + buf_size_bytes, read_buffer_size_bytes_); + read_buffer_ = buf; + read_buffer_size_bytes_ = buf_size_bytes; + } + + int64_t GetWriteBuffer(uint8_t* buf, size_t read_len_bytes) { + FARF(ALWAYS, "HexagonIOHandler GetWriteBuffer called, read_len_bytes: %d", read_len_bytes); + return write_buffer_.sgetn(reinterpret_cast(buf), read_len_bytes); + } + + void Close() { FARF(ALWAYS, "HexagonIOHandler Close called"); } + + void Exit(int code) { exit(code); } + + private: + const uint8_t* read_buffer_; + uint32_t read_buffer_size_bytes_; + + std::stringbuf write_buffer_; +}; + +class HexagonRPCServer { + public: + explicit HexagonRPCServer(uint8_t* receive_buffer) : io_{receive_buffer}, rpc_server_{&io_} {}; + + int64_t Write(const uint8_t* data, size_t data_len_bytes) { + io_.SetReadBuffer(data, data_len_bytes); + rpc_server_.ProcessOnePacket(); + return (int64_t)data_len_bytes; + } + + int64_t Read(uint8_t* buf, size_t read_len_bytes) { + return io_.GetWriteBuffer(buf, read_len_bytes); + } + + private: + HexagonIOHandler io_; + MinRPCServer rpc_server_; +}; + +} // namespace hexagon +} // namespace runtime +} // namespace tvm + +static tvm::runtime::hexagon::HexagonRPCServer* g_hexagon_rpc_server = nullptr; + +static AEEResult hexagon_rpc_server_init() { + uint8_t* receive_buffer = new uint8_t[TVM_HEXAGON_RPC_BUFF_SIZE_BYTES]; + tvm::runtime::hexagon::HexagonRPCServer* rpc_server = + new tvm::runtime::hexagon::HexagonRPCServer(receive_buffer); + g_hexagon_rpc_server = rpc_server; +} + +const tvm::runtime::PackedFunc get_runtime_func(const std::string& name) { + if (const tvm::runtime::PackedFunc* pf = tvm::runtime::Registry::Get(name)) { + return *pf; + } + return tvm::runtime::PackedFunc(); +} + +void reset_device_api() { + const tvm::runtime::PackedFunc api = get_runtime_func("device_api.hexagon.v2"); + tvm::runtime::Registry::Register("device_api.hexagon", true).set_body(api); +} + +int __QAIC_HEADER(hexagon_rpc_open)(const char* uri, remote_handle64* handle) { + *handle = static_cast(reinterpret_cast(malloc(1))); + if (!*handle) { + FARF(ERROR, "%s: cannot allocate memory", __func__); + return AEE_ENOMEMORY; + } + reset_device_api(); + hexagon_rpc_server_init(); + return AEE_SUCCESS; +} + +int __QAIC_HEADER(hexagon_rpc_close)(remote_handle64 handle) { + FARF(ALWAYS, "%s", __func__); + if (handle) { + free(reinterpret_cast(static_cast(handle))); + } + return AEE_SUCCESS; +} + +// Send from Host to Hexagon over Android +AEEResult __QAIC_HEADER(hexagon_rpc_send)(remote_handle64 _handle, const unsigned char* data, + int dataLen) { + if (g_hexagon_rpc_server == nullptr) { + FARF(ERROR, "RPC Server is not initialized."); + return AEE_EFAILED; + } + + int64_t written_size = g_hexagon_rpc_server->Write(reinterpret_cast(data), + static_cast(dataLen)); + if (written_size != dataLen) { + FARF(ERROR, "RPC Server Write failed, written_size (%d) != dataLen (%d)", written_size, + dataLen); + return AEE_EFAILED; + } + return AEE_SUCCESS; +} + +// Receive from Hexagon and send to Host over Android. +AEEResult __QAIC_HEADER(hexagon_rpc_receive)(remote_handle64 _handle, unsigned char* data, + int dataLen, int64_t* buf_written_size) { + int64_t read_size = + g_hexagon_rpc_server->Read(reinterpret_cast(data), static_cast(dataLen)); + *buf_written_size = read_size; + if (read_size == dataLen) { + return AEE_SUCCESS; + } else { + FARF(ALWAYS, "RPC Server Read failed, read_size (%d) != dataLen (%d)", read_size, dataLen); + return AEE_EFAILED; + } +} + +TVM_REGISTER_GLOBAL("tvm.hexagon.load_module") + .set_body([](tvm::runtime::TVMArgs args, tvm::runtime::TVMRetValue* rv) { + std::string soname = args[0]; + tvm::ObjectPtr n = tvm::runtime::CreateDSOLibraryObject(soname); + *rv = CreateModuleFromLibrary(n, tvm::runtime::hexagon::WrapPackedFunc); + }); diff --git a/src/runtime/hexagon/rpc/hexagon_rpc.idl b/src/runtime/hexagon/rpc/hexagon_rpc.idl new file mode 100644 index 000000000000..858ca501cf58 --- /dev/null +++ b/src/runtime/hexagon/rpc/hexagon_rpc.idl @@ -0,0 +1,28 @@ +/* + * Licensed to the Apache Software Foundation (ASF) under one + * or more contributor license agreements. See the NOTICE file + * distributed with this work for additional information + * regarding copyright ownership. The ASF licenses this file + * to you under the Apache License, Version 2.0 (the + * "License"); you may not use this file except in compliance + * with the License. You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, + * software distributed under the License is distributed on an + * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY + * KIND, either express or implied. See the License for the + * specific language governing permissions and limitations + * under the License. + */ + +#include "remote.idl" +#include "AEEStdDef.idl" + +typedef sequence buffer; + +interface hexagon_rpc : remote_handle64 { + AEEResult send(in buffer buf); + AEEResult receive(rout buffer buf, rout int64_t buf_written_size); +}; diff --git a/src/target/llvm/codegen_hexagon.cc b/src/target/llvm/codegen_hexagon.cc index 0cc5c5bfe83a..fed2ad5937ee 100644 --- a/src/target/llvm/codegen_hexagon.cc +++ b/src/target/llvm/codegen_hexagon.cc @@ -836,9 +836,9 @@ runtime::Module BuildHexagon(IRModule mod, Target target) { std::string so_name(o_name, 0, o_name.size() - 1); so_name += "so"; - const auto* f = tvm::runtime::Registry::Get("tvm.contrib.hexagon.link_shared"); - ICHECK(f != nullptr) << "tvm.contrib.hexagon.link_shared does not to exist, " - "do import tvm.contrib.hexagon"; + const auto* f = tvm::runtime::Registry::Get("tvm.contrib.hexagon.hexagon.link_shared"); + ICHECK(f != nullptr) << "tvm.contrib.hexagon.hexagon.link_shared does not to exist, " + "do import tvm.contrib.hexagon.hexagon"; Array o_names = {StringImm(o_name)}; int rc = (*f)(so_name, o_names); diff --git a/src/target/llvm/llvm_common.h b/src/target/llvm/llvm_common.h index fcc44fb8f95c..f31f3f698ddb 100644 --- a/src/target/llvm/llvm_common.h +++ b/src/target/llvm/llvm_common.h @@ -27,7 +27,6 @@ #ifdef _MSC_VER #pragma warning(disable : 4141 4291 4146 4624) #endif - #ifdef TVM_LLVM_VERSION #include @@ -72,11 +71,7 @@ #include #include #include -#if TVM_LLVM_VERSION >= 140 -#include -#else #include -#endif #include #include #include diff --git a/tests/lint/check_file_type.py b/tests/lint/check_file_type.py index 1b45ac783c29..964003845961 100644 --- a/tests/lint/check_file_type.py +++ b/tests/lint/check_file_type.py @@ -150,6 +150,8 @@ "apps/microtvm/reference-vm/arduino/base-box/Vagrantfile.packer-template", "apps/microtvm/reference-vm/zephyr/Vagrantfile", "apps/microtvm/reference-vm/zephyr/base-box/Vagrantfile.packer-template", + # Hexagon + "src/runtime/hexagon/rpc/android_bash.sh.template", } diff --git a/tests/python/contrib/test_hexagon/proxy_rpc/test_matmul.py b/tests/python/contrib/test_hexagon/proxy_rpc/test_matmul.py index 2b18911bacf2..b83ce4480445 100644 --- a/tests/python/contrib/test_hexagon/proxy_rpc/test_matmul.py +++ b/tests/python/contrib/test_hexagon/proxy_rpc/test_matmul.py @@ -21,7 +21,7 @@ import tvm import tvm.testing from tvm import te -import tvm.contrib.hexagon as hexagon +import tvm.contrib.hexagon.hexagon as hexagon from tvm.contrib import utils from tvm import rpc import numpy as np diff --git a/tests/python/contrib/test_hexagon/rpc/conftest.py b/tests/python/contrib/test_hexagon/rpc/conftest.py new file mode 100644 index 000000000000..39806b29bd96 --- /dev/null +++ b/tests/python/contrib/test_hexagon/rpc/conftest.py @@ -0,0 +1,82 @@ +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you under the Apache License, Version 2.0 (the +# "License"); you may not use this file except in compliance +# with the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, +# software distributed under the License is distributed on an +# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +# KIND, either express or implied. See the License for the +# specific language governing permissions and limitations +# under the License. +import pytest +import os + +import tvm.testing + + +def pytest_addoption(parser): + parser.addoption( + "--serial-number", + required=True, + help=("Android device serial number list from 'adb' command."), + ) + + +@pytest.fixture +def android_serial_number(request): + return request.config.getoption("--serial-number") + + +@tvm.testing.fixture +def tvm_tracker_host(): + return os.environ["TVM_TRACKER_HOST"] + + +@tvm.testing.fixture +def tvm_tracker_port(): + return int(os.environ["TVM_TRACKER_PORT"]) + + +def _compose(args, decs): + """Helper to apply multiple markers""" + if len(args) > 0: + f = args[0] + for d in reversed(decs): + f = d(f) + return f + return decs + + +def requires_rpc_tracker(*args): + """Mark a test as requiring an RPC tracker to exist in + the host environment to run.""" + _requires_rpc_tracker = [ + *tvm.testing.requires_rpc(), + pytest.mark.skipif( + os.environ.get("TVM_TRACKER_HOST") == None, + reason="Missing environment variable, TVM_TRACKER_HOST", + ), + pytest.mark.skipif( + os.environ.get("TVM_TRACKER_PORT") == None, + reason="Missing environment variable, TVM_TRACKER_PORT", + ), + ] + + return _compose(args, _requires_rpc_tracker) + + +def requires_ndk_cc(*args): + _requires_ndk_cc = [ + pytest.mark.skipif( + os.environ.get("TVM_NDK_CC") == None, + reason="TVM_NDK_CC environment variable is required to run this test.", + ), + ] + + return _compose(args, _requires_ndk_cc) diff --git a/tests/python/contrib/test_hexagon/rpc/test_launcher.md b/tests/python/contrib/test_hexagon/rpc/test_launcher.md new file mode 100644 index 000000000000..47daf0928455 --- /dev/null +++ b/tests/python/contrib/test_hexagon/rpc/test_launcher.md @@ -0,0 +1,98 @@ + + + + + + + + + + + + + + + + + +# HexagonLauncher +HexagonLauncher is a class to handle interactions with an Android phone which includes Hexagon DSP to run a TVMModule(function/operation/graph) on Hexagon. HexagonLauncher reuses minRPC implementation to setup an RPC connection from host (your local machine) to Hexagon target which is passed through Android RPC server. + +## Build Required Tools/Libraries +Here are the steps that are taken to prepare a runtime on a Hexagon device to test any model. + +- Build TVMRuntime library and C++ RPC server for Android. +- Build minRPC server along with FastRPC for Hexagon. +- Build TVM library with Hexagon support for host machine. +- Build TVMRuntime library and C++ RPC server for host machine. + +To build these pieces, you can use a cmake command as follow. + +```bash +cmake -DUSE_HEXAGON_RPC=ON \ + -DUSE_ANDROID_TOOLCHAIN=/path/to/android-ndk/build/cmake/android.toolchain.cmake \ + -DANDROID_PLATFORM=android-28 \ + -DANDROID_ABI=arm64-v8a \ + -DUSE_HEXAGON_ARCH=v65|v66|v68 \ + -DUSE_HEXAGON_SDK=/path/to/Hexagon/SDK \ + -DUSE_HEXAGON_TOOLCHAIN=/path/to/Hexagon/toolchain/ \ + -DUSE_LLVM=/path/to/llvm/bin/llvm-config \ + -DUSE_CPP_RPC=ON \ + -DCMAKE_CXX_COMPILER=/path/to/clang++ \ + -DCMAKE_CXX_FLAGS='-stdlib=libc++' .. +``` + +## Testing Using HexagonLauncher +Before starting a test you need to run an RPC tracker on your local machine and export HOST and PORT as environment variables + +```bash +export TVM_TRACKER_HOST="0.0.0.0" +export TVM_TRACKER_PORT=9192 +python -m tvm.exec.rpc_tracker --host $TVM_TRACKER_HOST --port $TVM_TRACKER_PORT +``` + +Now, follow these steps to create an RPC session from host to Hexagon. + +```python +# create an HexagonLauncher instance +launcher = HexagonLauncher(serial_number="Serial number taken from `adb devices` command") + +# Create a workspace directory for this test on Android. +# Upload required Android artifacts including TVMRuntime library and RPC server to Android workspace. +# Uses port `forward` and `reverse` to open connection on certain ports that TVM uses to connect to RPC tracker. +# Execute `android_bash.sh` on Android which creates two RPC servers and connects them to RPC tracker running on host machine. +launcher.android_run_rpc(rpc_tracker_host="TVM_TRACKER_HOST", rpc_tracker_port="TVM_TRACKER_PORT") + +# Upload Hexagon RPC libraries to Android workspace. +launcher.hexagon_setup() + +# Connect and RPC session from host to Android to be able to upload TVMModules to Android file system. +remote_kw = { + "host": "TVM_TRACKER_HOST", + "port": "TVM_TRACKER_PORT", + "priority": 0, + "timeout": 60, +} +launcher.android_remote_setup(remote_kw) + +# Create an RPC session from host to Hexagon. +launcher.hexagon_session_setup(remote_kw) + +# Upload TVMModule binary file to Android remote. +launcher.upload("Path to DSO binary file on host", "DSO filename on Android remote") +``` + +- To execute a single function/operator on Hexagon, follow these steps. + ```python + # Enter session. + with launcher.session as sess: + # dlopen DSO binary file on Hexagon. + mod = launcher.get_module(dso_binary) + # Use mod to run function/operator on Hexagon... + ``` +- Or, follow these steps to create a GraphExecutor and run a JSON graph. + ```python + graph_mod = launcher.get_local_graph_executor(lowered, dso_binary) + graph_mod.set_input(...) + graph_mod.run(...) + ``` diff --git a/tests/python/contrib/test_hexagon/rpc/test_launcher.py b/tests/python/contrib/test_hexagon/rpc/test_launcher.py new file mode 100644 index 000000000000..00f18a961ea4 --- /dev/null +++ b/tests/python/contrib/test_hexagon/rpc/test_launcher.py @@ -0,0 +1,260 @@ +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you under the Apache License, Version 2.0 (the +# "License"); you may not use this file except in compliance +# with the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, +# software distributed under the License is distributed on an +# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +# KIND, either express or implied. See the License for the +# specific language governing permissions and limitations +# under the License. + +import sys +import pytest +import numpy as np +import os + +import tvm.testing +from tvm import te +from tvm import relay +from tvm.relay.backend import Executor, Runtime +from tvm.contrib import utils, ndk +from tvm.contrib.hexagon.build import HexagonLauncher +import tvm.contrib.hexagon.hexagon as hexagon + +from conftest import requires_rpc_tracker, requires_ndk_cc + + +@requires_rpc_tracker +@requires_ndk_cc +def test_rpc_on_android(tvm_tracker_host, tvm_tracker_port, android_serial_number): + if "TVM_NDK_CC" not in os.environ: + raise RuntimeError( + "Require environment variable TVM_NDK_CC" " to be the NDK standalone compiler" + ) + target = "llvm -mtriple=arm64-linux-android" + + n = tvm.runtime.convert(1024) + A = te.placeholder((n,), name="A") + B = te.compute(A.shape, lambda *i: A(*i) + 1.0, name="B") + a_np = np.random.uniform(size=1024).astype(A.dtype) + temp = utils.tempdir() + + s = te.create_schedule(B.op) + xo, xi = s[B].split(B.op.axis[0], factor=64) + s[B].parallel(xi) + s[B].pragma(xo, "parallel_launch_point") + s[B].pragma(xi, "parallel_barrier_when_finish") + f = tvm.build(s, [A, B], target, name="myadd_cpu") + path_dso_cpu = temp.relpath("cpu_lib.so") + f.export_library(path_dso_cpu, ndk.create_shared) + + launcher = HexagonLauncher(serial_number=android_serial_number) + launcher.android_run_rpc(rpc_tracker_host=tvm_tracker_host, rpc_tracker_port=tvm_tracker_port) + + remote_kw = { + "host": tvm_tracker_host, + "port": tvm_tracker_port, + "priority": 0, + "timeout": 60, + } + launcher.android_remote_setup(remote_kw) + + print("Run CPU test ...") + remote = launcher.android_remote + dev = remote.cpu(0) + launcher.upload(path_dso_cpu) + f2 = remote.load_module("cpu_lib.so") + a = tvm.nd.array(a_np, dev) + b = tvm.nd.array(np.zeros(1024, dtype=A.dtype), dev) + time_f = f2.time_evaluator(f2.entry_name, dev, number=10) + cost = time_f(a, b).mean + print("%g secs/op\n" % cost) + np.testing.assert_equal(b.numpy(), a.numpy() + 1) + launcher.close() + + +@requires_rpc_tracker +def test_add(tvm_tracker_host, tvm_tracker_port, android_serial_number): + dtype = "int8" + A = tvm.te.placeholder((2,), dtype=dtype) + B = tvm.te.placeholder((1,), dtype=dtype) + C = tvm.te.compute(A.shape, lambda i: A[i] + B[0], name="C") + sched = tvm.te.create_schedule(C.op) + + target = tvm.target.hexagon("v68", link_params=True) + func = tvm.build(sched, [A, B, C], target=target, target_host=target, name="add") + + temp = utils.tempdir() + dso_binary = "test_binary.so" + dso_binary_path = temp.relpath(dso_binary) + func.save(dso_binary_path) + + launcher = HexagonLauncher(serial_number=android_serial_number) + launcher.android_run_rpc(rpc_tracker_host=tvm_tracker_host, rpc_tracker_port=tvm_tracker_port) + launcher.hexagon_setup() + remote_kw = { + "host": tvm_tracker_host, + "port": tvm_tracker_port, + "priority": 0, + "timeout": 60, + } + launcher.android_remote_setup(remote_kw) + launcher.hexagon_session_setup(remote_kw) + launcher.upload(dso_binary_path, dso_binary) + + with launcher.session as sess: + mod = launcher.get_module(dso_binary) + A_data = tvm.nd.array(np.array([2, 3], dtype=dtype), device=sess.device) + assert (A_data.numpy() == np.array([2, 3])).all() + B_data = tvm.nd.array(np.array([4], dtype=dtype), device=sess.device) + assert (B_data.numpy() == np.array([4])).all() + C_data = tvm.nd.array(np.array([0, 0], dtype=dtype), device=sess.device) + assert (C_data.numpy() == np.array([0, 0])).all() + + mod["add"](A_data, B_data, C_data) + assert (C_data.numpy() == np.array([6, 7])).all() + launcher.close() + + +class TestMatMul: + M = tvm.testing.parameter(32) + N = tvm.testing.parameter(32) + K = tvm.testing.parameter(32) + + @requires_rpc_tracker + def test_matmul(self, tvm_tracker_host, tvm_tracker_port, android_serial_number, M, N, K): + X = te.placeholder((M, K), dtype="float32") + Y = te.placeholder((K, N), dtype="float32") + k1 = te.reduce_axis((0, K), name="k1") + Z = te.compute((M, N), lambda i, j: te.sum(X[i, k1] * Y[k1, j], axis=[k1])) + schedule = te.create_schedule(Z.op) + + target_hexagon = tvm.target.hexagon("v68", link_params=True) + func = tvm.build(schedule, [X, Y, Z], target=target_hexagon, target_host=target_hexagon) + + temp = utils.tempdir() + dso_binary = "test_binary.so" + dso_binary_path = temp.relpath(dso_binary) + func.save(dso_binary_path) + + launcher = HexagonLauncher(serial_number=android_serial_number) + launcher.android_run_rpc( + rpc_tracker_host=tvm_tracker_host, rpc_tracker_port=tvm_tracker_port + ) + launcher.hexagon_setup() + remote_kw = { + "host": tvm_tracker_host, + "port": tvm_tracker_port, + "priority": 0, + "timeout": 60, + } + launcher.android_remote_setup(remote_kw) + launcher.hexagon_session_setup(remote_kw) + launcher.upload(dso_binary_path, dso_binary) + + x = np.random.uniform(size=[i.value for i in X.shape]).astype(X.dtype) + y = np.random.uniform(size=[i.value for i in Y.shape]).astype(Y.dtype) + z = np.zeros([i.value for i in Z.shape], dtype=Z.dtype) + + with launcher.session as sess: + mod = launcher.get_module(dso_binary) + xt = tvm.nd.array(x, device=sess.device) + yt = tvm.nd.array(y, device=sess.device) + zt = tvm.nd.array(z, device=sess.device) + mod(xt, yt, zt) + + target_llvm = tvm.target.Target("llvm") + mod = tvm.build(schedule, [X, Y, Z], target=target_llvm, target_host=target_llvm) + device = tvm.cpu(0) + xtcpu = tvm.nd.array(x, device) + ytcpu = tvm.nd.array(y, device) + ztcpu = tvm.nd.array(z, device) + mod(xtcpu, ytcpu, ztcpu) + launcher.close() + + tvm.testing.assert_allclose(zt.asnumpy(), ztcpu.asnumpy(), rtol=1e-4) + + +@requires_rpc_tracker +def test_graph_executor(tvm_tracker_host, tvm_tracker_port, android_serial_number): + dtype = "float32" + data = relay.var("data", relay.TensorType((1, 64, 64, 3), dtype)) + weight = relay.var("weight", relay.TensorType((5, 5, 3, 8), dtype)) + y = relay.nn.conv2d( + data, + weight, + padding=(2, 2), + kernel_size=(5, 5), + data_layout="NHWC", + kernel_layout="HWIO", + out_dtype="float32", + ) + f = relay.Function([data, weight], y) + relay_mod = tvm.IRModule.from_expr(f) + relay_mod = relay.transform.InferType()(relay_mod) + + target_hexagon = tvm.target.hexagon("v68") + runtime = Runtime("cpp") + executor = Executor("graph") + + temp = utils.tempdir() + dso_binary = "test_binary.so" + dso_binary_path = temp.relpath(dso_binary) + + with tvm.transform.PassContext(opt_level=3): + lowered = tvm.relay.build( + relay_mod, + target=target_hexagon, + target_host=target_hexagon, + runtime=runtime, + executor=executor, + ) + lowered.get_lib().save(dso_binary_path) + + launcher = HexagonLauncher(serial_number=android_serial_number) + launcher.android_run_rpc(rpc_tracker_host=tvm_tracker_host, rpc_tracker_port=tvm_tracker_port) + launcher.hexagon_setup() + remote_kw = { + "host": tvm_tracker_host, + "port": tvm_tracker_port, + "priority": 0, + "timeout": 60, + } + launcher.android_remote_setup(remote_kw) + launcher.hexagon_session_setup(remote_kw) + launcher.upload(dso_binary_path, dso_binary) + + graph_mod = launcher.get_local_graph_executor(lowered, dso_binary) + weight_in = np.random.rand(5, 5, 3, 8).astype(dtype=dtype) + data_in = np.random.rand(1, 64, 64, 3).astype(dtype=dtype) + graph_mod.set_input(weight=weight_in) + graph_mod.run(data=data_in) + hexagon_output = graph_mod.get_output(0).numpy() + + target_llvm = tvm.target.Target("llvm") + with tvm.transform.PassContext(opt_level=3): + llvm_lowered = tvm.relay.build( + relay_mod, + target=target_llvm, + target_host=target_llvm, + runtime=runtime, + executor=executor, + ) + llvm_graph_mod = tvm.contrib.graph_executor.GraphModule(llvm_lowered["default"](tvm.cpu(0))) + llvm_graph_mod.set_input(weight=weight_in) + llvm_graph_mod.run(data=data_in) + expected_output = llvm_graph_mod.get_output(0).numpy() + launcher.close() + + tvm.testing.assert_allclose(hexagon_output, expected_output, rtol=1e-4, atol=1e-5) + + +if __name__ == "__main__": + sys.exit(pytest.main(sys.argv)) diff --git a/tests/python/unittest/test_target_codegen_hexagon.py b/tests/python/unittest/test_target_codegen_hexagon.py index 220fa869d840..ef0eb4ff5a7e 100644 --- a/tests/python/unittest/test_target_codegen_hexagon.py +++ b/tests/python/unittest/test_target_codegen_hexagon.py @@ -23,12 +23,12 @@ import tvm import tvm.relay import tvm.testing -import tvm.contrib.hexagon as hexagon +import tvm.contrib.hexagon.hexagon as hexagon @pytest.fixture(autouse=True) def register_linker(): - original_linker = tvm.contrib.hexagon.hexagon_link() + original_linker = tvm.contrib.hexagon.hexagon.hexagon_link() # Register a phony linker, so that we can test codegen without a Hexagon toolchain. hexagon.register_linker(lambda: "/bin/true") yield None From 87d8bec172a1fb9feaf561bfab6e3d0afa5f125f Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Fri, 3 Dec 2021 13:49:32 -0800 Subject: [PATCH 02/15] removed android remote and updated Readme --- python/tvm/contrib/hexagon/build.py | 46 ++++------- .../hexagon/rpc/android_bash.sh.template | 11 +-- .../contrib/test_hexagon/rpc/conftest.py | 10 +-- .../contrib/test_hexagon/rpc/test_launcher.md | 10 +-- .../contrib/test_hexagon/rpc/test_launcher.py | 77 ++++--------------- 5 files changed, 41 insertions(+), 113 deletions(-) diff --git a/python/tvm/contrib/hexagon/build.py b/python/tvm/contrib/hexagon/build.py index 393e60b43594..191cc1c3ad6b 100644 --- a/python/tvm/contrib/hexagon/build.py +++ b/python/tvm/contrib/hexagon/build.py @@ -25,7 +25,6 @@ import datetime import tvm -from tvm import rpc as _rpc from ..._ffi import libinfo from .session import Session @@ -75,12 +74,10 @@ def __init__(self, serial_number: str): self._serial_number = serial_number self._adb_device_sub_cmd = ["adb", "-s", self._serial_number] - self._android_remote = None self._mod = None self._workspace = None - ANDROID_REMOTE_DEVICE_KEY = "hexagon_dev-android" - HEXAGON_REMOTE_DEVICE_KEY = "hexagon_dev-hexagon" + HEXAGON_REMOTE_DEVICE_KEY = "hexagon-dev" def android_run_rpc( self, @@ -127,10 +124,6 @@ def android_run_rpc( line = line.replace("", str(rpc_tracker_host)) if "" in line: line = line.replace("", str(rpc_tracker_port)) - if f"" in line: - line = line.replace( - "", self.ANDROID_REMOTE_DEVICE_KEY - ) if "" in line: line = line.replace( "", self.HEXAGON_REMOTE_DEVICE_KEY @@ -184,21 +177,6 @@ def hexagon_setup(self): dst_path = f"{self._workspace}/{item}" subprocess.check_call(self._adb_device_sub_cmd + ["push", src_path, dst_path]) - def android_remote_setup(self, remote_kw: dict): - """Setup remote connection from host to Android RPC server. - - Parameters - ---------- - remote_kw : dict - RPC tracker configs. - """ - tracker = _rpc.connect_tracker(remote_kw["host"], remote_kw["port"]) - self.android_remote = tracker.request( - self.ANDROID_REMOTE_DEVICE_KEY, - priority=remote_kw["priority"], - session_timeout=remote_kw["timeout"], - ) - def hexagon_session_setup(self, remote_kw: dict): """Setup Hexagon RPC Session from host to Hexagon device. @@ -207,7 +185,6 @@ def hexagon_session_setup(self, remote_kw: dict): remote_kw : dict RPC tracker configs. """ - assert self.android_remote, "android_remote should initialied first." hexagon_remote_kw = dict(remote_kw) hexagon_remote_kw["key"] = self.HEXAGON_REMOTE_DEVICE_KEY self.session = Session(hexagon_remote_kw) @@ -229,7 +206,7 @@ def get_module(self, module_name: str): self._mod = self.session.load_module(module_path) return self._mod - def upload(self, host_path: Union[str, pathlib.Path], remote_filename: str = None): + def upload(self, host_path: Union[str, pathlib.Path], remote_filename: str): """Upload a file to remote(Android). Parameters @@ -245,11 +222,8 @@ def upload(self, host_path: Union[str, pathlib.Path], remote_filename: str = Non A TVM Module loaded on hexagon. """ src_path = str(host_path) - if remote_filename: - dst_remote_path = f"{self._workspace}/{remote_filename}" - else: - dst_remote_path = None - self.android_remote.upload(src_path, target=dst_remote_path) + dst_remote_path = f"{self._workspace}/{remote_filename}" + subprocess.check_call(self._adb_device_sub_cmd + ["push", src_path, dst_remote_path]) def get_local_graph_executor(self, libmod, remote_libmod_filename: str): """Create a GraphModule. @@ -274,6 +248,12 @@ def get_local_graph_executor(self, libmod, remote_libmod_filename: str): ) def close(self): - """Close RPC servers""" - subprocess.Popen(self._adb_device_sub_cmd + ["shell", "kill `cat android_rpc_pid.txt`"]) - subprocess.Popen(self._adb_device_sub_cmd + ["shell", "kill `cat hexagon_rpc_pid.txt`"]) + """Close RPC server on Android""" + # kill process childs + subprocess.Popen( + self._adb_device_sub_cmd + ["shell", f"pkill -P `cat {self._workspace}/rpc_pid.txt`"] + ) + # kill main process + subprocess.Popen( + self._adb_device_sub_cmd + ["shell", f"kill `cat {self._workspace}/rpc_pid.txt`"] + ) diff --git a/src/runtime/hexagon/rpc/android_bash.sh.template b/src/runtime/hexagon/rpc/android_bash.sh.template index f6afa793147c..7bf6d773f2f7 100644 --- a/src/runtime/hexagon/rpc/android_bash.sh.template +++ b/src/runtime/hexagon/rpc/android_bash.sh.template @@ -17,13 +17,8 @@ # under the License. export LD_LIBRARY_PATH=. -./tvm_rpc_android server --port= --tracker=: --key=& -android_rpc_pid=$! ./tvm_rpc_android server --port= --tracker=: --key=& -hexagon_rpc_pid=$! +rpc_pid=$! -rm -f android_rpc_pid.txt -echo $android_rpc_pid >> android_rpc_pid.txt - -rm -f hexagon_rpc_pid.txt -echo $hexagon_rpc_pid >> hexagon_rpc_pid.txt +rm -f rpc_pid.txt +echo $rpc_pid >> rpc_pid.txt diff --git a/tests/python/contrib/test_hexagon/rpc/conftest.py b/tests/python/contrib/test_hexagon/rpc/conftest.py index 39806b29bd96..24edd0195504 100644 --- a/tests/python/contrib/test_hexagon/rpc/conftest.py +++ b/tests/python/contrib/test_hexagon/rpc/conftest.py @@ -71,12 +71,12 @@ def requires_rpc_tracker(*args): return _compose(args, _requires_rpc_tracker) -def requires_ndk_cc(*args): - _requires_ndk_cc = [ +def requires_hexagon_toolchain(*args): + _requires_hexagon_toolchain = [ pytest.mark.skipif( - os.environ.get("TVM_NDK_CC") == None, - reason="TVM_NDK_CC environment variable is required to run this test.", + os.environ.get("HEXAGON_TOOLCHAIN") == None, + reason="HEXAGON_TOOLCHAIN environment variable is required to run this test.", ), ] - return _compose(args, _requires_ndk_cc) + return _compose(args, _requires_hexagon_toolchain) diff --git a/tests/python/contrib/test_hexagon/rpc/test_launcher.md b/tests/python/contrib/test_hexagon/rpc/test_launcher.md index 47daf0928455..463b88e3f374 100644 --- a/tests/python/contrib/test_hexagon/rpc/test_launcher.md +++ b/tests/python/contrib/test_hexagon/rpc/test_launcher.md @@ -43,9 +43,12 @@ cmake -DUSE_HEXAGON_RPC=ON \ ``` ## Testing Using HexagonLauncher -Before starting a test you need to run an RPC tracker on your local machine and export HOST and PORT as environment variables +Before starting a test you need to run an RPC tracker on your local machine and export HOST and PORT as environment variables. Also, you need to export Clang libraries to `LD_LIBRARY_PATH` and Hexagon toolchain to `HEXAGON_TOOLCHAIN`. ```bash +export LD_LIBRARY_PATH="$LD_LIBRARY_PATH:/path/to/clang++/lib" +export HEXAGON_TOOLCHAIN="/path/to/Hexagon/toolchain/" + export TVM_TRACKER_HOST="0.0.0.0" export TVM_TRACKER_PORT=9192 python -m tvm.exec.rpc_tracker --host $TVM_TRACKER_HOST --port $TVM_TRACKER_PORT @@ -66,16 +69,13 @@ launcher.android_run_rpc(rpc_tracker_host="TVM_TRACKER_HOST", rpc_tracker_port=" # Upload Hexagon RPC libraries to Android workspace. launcher.hexagon_setup() -# Connect and RPC session from host to Android to be able to upload TVMModules to Android file system. +# Create an RPC session from host to Hexagon. remote_kw = { "host": "TVM_TRACKER_HOST", "port": "TVM_TRACKER_PORT", "priority": 0, "timeout": 60, } -launcher.android_remote_setup(remote_kw) - -# Create an RPC session from host to Hexagon. launcher.hexagon_session_setup(remote_kw) # Upload TVMModule binary file to Android remote. diff --git a/tests/python/contrib/test_hexagon/rpc/test_launcher.py b/tests/python/contrib/test_hexagon/rpc/test_launcher.py index 00f18a961ea4..1a1190b8f769 100644 --- a/tests/python/contrib/test_hexagon/rpc/test_launcher.py +++ b/tests/python/contrib/test_hexagon/rpc/test_launcher.py @@ -28,59 +28,11 @@ from tvm.contrib.hexagon.build import HexagonLauncher import tvm.contrib.hexagon.hexagon as hexagon -from conftest import requires_rpc_tracker, requires_ndk_cc - - -@requires_rpc_tracker -@requires_ndk_cc -def test_rpc_on_android(tvm_tracker_host, tvm_tracker_port, android_serial_number): - if "TVM_NDK_CC" not in os.environ: - raise RuntimeError( - "Require environment variable TVM_NDK_CC" " to be the NDK standalone compiler" - ) - target = "llvm -mtriple=arm64-linux-android" - - n = tvm.runtime.convert(1024) - A = te.placeholder((n,), name="A") - B = te.compute(A.shape, lambda *i: A(*i) + 1.0, name="B") - a_np = np.random.uniform(size=1024).astype(A.dtype) - temp = utils.tempdir() - - s = te.create_schedule(B.op) - xo, xi = s[B].split(B.op.axis[0], factor=64) - s[B].parallel(xi) - s[B].pragma(xo, "parallel_launch_point") - s[B].pragma(xi, "parallel_barrier_when_finish") - f = tvm.build(s, [A, B], target, name="myadd_cpu") - path_dso_cpu = temp.relpath("cpu_lib.so") - f.export_library(path_dso_cpu, ndk.create_shared) - - launcher = HexagonLauncher(serial_number=android_serial_number) - launcher.android_run_rpc(rpc_tracker_host=tvm_tracker_host, rpc_tracker_port=tvm_tracker_port) - - remote_kw = { - "host": tvm_tracker_host, - "port": tvm_tracker_port, - "priority": 0, - "timeout": 60, - } - launcher.android_remote_setup(remote_kw) - - print("Run CPU test ...") - remote = launcher.android_remote - dev = remote.cpu(0) - launcher.upload(path_dso_cpu) - f2 = remote.load_module("cpu_lib.so") - a = tvm.nd.array(a_np, dev) - b = tvm.nd.array(np.zeros(1024, dtype=A.dtype), dev) - time_f = f2.time_evaluator(f2.entry_name, dev, number=10) - cost = time_f(a, b).mean - print("%g secs/op\n" % cost) - np.testing.assert_equal(b.numpy(), a.numpy() + 1) - launcher.close() +from conftest import requires_rpc_tracker, requires_hexagon_toolchain @requires_rpc_tracker +@requires_hexagon_toolchain def test_add(tvm_tracker_host, tvm_tracker_port, android_serial_number): dtype = "int8" A = tvm.te.placeholder((2,), dtype=dtype) @@ -88,8 +40,10 @@ def test_add(tvm_tracker_host, tvm_tracker_port, android_serial_number): C = tvm.te.compute(A.shape, lambda i: A[i] + B[0], name="C") sched = tvm.te.create_schedule(C.op) - target = tvm.target.hexagon("v68", link_params=True) - func = tvm.build(sched, [A, B, C], target=target, target_host=target, name="add") + target_hexagon = tvm.target.hexagon("v68", link_params=True) + func = tvm.build( + sched, [A, B, C], tvm.target.Target(target_hexagon, host=target_hexagon), name="add" + ) temp = utils.tempdir() dso_binary = "test_binary.so" @@ -105,7 +59,6 @@ def test_add(tvm_tracker_host, tvm_tracker_port, android_serial_number): "priority": 0, "timeout": 60, } - launcher.android_remote_setup(remote_kw) launcher.hexagon_session_setup(remote_kw) launcher.upload(dso_binary_path, dso_binary) @@ -129,6 +82,7 @@ class TestMatMul: K = tvm.testing.parameter(32) @requires_rpc_tracker + @requires_hexagon_toolchain def test_matmul(self, tvm_tracker_host, tvm_tracker_port, android_serial_number, M, N, K): X = te.placeholder((M, K), dtype="float32") Y = te.placeholder((K, N), dtype="float32") @@ -137,7 +91,9 @@ def test_matmul(self, tvm_tracker_host, tvm_tracker_port, android_serial_number, schedule = te.create_schedule(Z.op) target_hexagon = tvm.target.hexagon("v68", link_params=True) - func = tvm.build(schedule, [X, Y, Z], target=target_hexagon, target_host=target_hexagon) + func = tvm.build( + schedule, [X, Y, Z], tvm.target.Target(target_hexagon, host=target_hexagon) + ) temp = utils.tempdir() dso_binary = "test_binary.so" @@ -155,7 +111,6 @@ def test_matmul(self, tvm_tracker_host, tvm_tracker_port, android_serial_number, "priority": 0, "timeout": 60, } - launcher.android_remote_setup(remote_kw) launcher.hexagon_session_setup(remote_kw) launcher.upload(dso_binary_path, dso_binary) @@ -171,7 +126,7 @@ def test_matmul(self, tvm_tracker_host, tvm_tracker_port, android_serial_number, mod(xt, yt, zt) target_llvm = tvm.target.Target("llvm") - mod = tvm.build(schedule, [X, Y, Z], target=target_llvm, target_host=target_llvm) + mod = tvm.build(schedule, [X, Y, Z], tvm.target.Target(target_llvm, host=target_llvm)) device = tvm.cpu(0) xtcpu = tvm.nd.array(x, device) ytcpu = tvm.nd.array(y, device) @@ -179,10 +134,11 @@ def test_matmul(self, tvm_tracker_host, tvm_tracker_port, android_serial_number, mod(xtcpu, ytcpu, ztcpu) launcher.close() - tvm.testing.assert_allclose(zt.asnumpy(), ztcpu.asnumpy(), rtol=1e-4) + tvm.testing.assert_allclose(zt.numpy(), ztcpu.numpy(), rtol=1e-4) @requires_rpc_tracker +@requires_hexagon_toolchain def test_graph_executor(tvm_tracker_host, tvm_tracker_port, android_serial_number): dtype = "float32" data = relay.var("data", relay.TensorType((1, 64, 64, 3), dtype)) @@ -211,8 +167,7 @@ def test_graph_executor(tvm_tracker_host, tvm_tracker_port, android_serial_numbe with tvm.transform.PassContext(opt_level=3): lowered = tvm.relay.build( relay_mod, - target=target_hexagon, - target_host=target_hexagon, + tvm.target.Target(target_hexagon, host=target_hexagon), runtime=runtime, executor=executor, ) @@ -227,7 +182,6 @@ def test_graph_executor(tvm_tracker_host, tvm_tracker_port, android_serial_numbe "priority": 0, "timeout": 60, } - launcher.android_remote_setup(remote_kw) launcher.hexagon_session_setup(remote_kw) launcher.upload(dso_binary_path, dso_binary) @@ -242,8 +196,7 @@ def test_graph_executor(tvm_tracker_host, tvm_tracker_port, android_serial_numbe with tvm.transform.PassContext(opt_level=3): llvm_lowered = tvm.relay.build( relay_mod, - target=target_llvm, - target_host=target_llvm, + tvm.target.Target(target_llvm, host=target_llvm), runtime=runtime, executor=executor, ) From cea8757ddc39a4c9589e1d5394a4ae431beb1660 Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Sat, 4 Dec 2021 02:16:25 -0800 Subject: [PATCH 03/15] Add check for workspace size --- python/tvm/contrib/hexagon/build.py | 22 +++++++++++++++++++++- 1 file changed, 21 insertions(+), 1 deletion(-) diff --git a/python/tvm/contrib/hexagon/build.py b/python/tvm/contrib/hexagon/build.py index 191cc1c3ad6b..3d8602c02b1a 100644 --- a/python/tvm/contrib/hexagon/build.py +++ b/python/tvm/contrib/hexagon/build.py @@ -61,7 +61,7 @@ def get_hexagon_rpc_dir() -> pathlib.Path: class HexagonLauncher: """Hexagon Launcher""" - def __init__(self, serial_number: str): + def __init__(self, serial_number: str, workspace_size_gb: int = 1): """Configure a new HexagonLauncher Parameters @@ -76,6 +76,7 @@ def __init__(self, serial_number: str): self._adb_device_sub_cmd = ["adb", "-s", self._serial_number] self._mod = None self._workspace = None + self._workspace_max_size_mb = workspace_size_gb * 1024 HEXAGON_REMOTE_DEVICE_KEY = "hexagon-dev" @@ -102,6 +103,10 @@ def android_run_rpc( rpc_tracker_port : int RPC tracker port on host """ + # Check size of base directory and cleanup if needed + while self._get_workspace_size() > self._workspace_max_size_mb: + self._workspace_remove_latest() + if not workspace_dir: self._workspace = str( ANDROID_HEXAGON_TEST_BASE_DIR @@ -257,3 +262,18 @@ def close(self): subprocess.Popen( self._adb_device_sub_cmd + ["shell", f"kill `cat {self._workspace}/rpc_pid.txt`"] ) + + def _get_workspace_size(self) -> int: + """Get workspace base directory size in MB""" + line = subprocess.check_output(self._adb_device_sub_cmd + ["shell", "du", "-shm", str(ANDROID_HEXAGON_TEST_BASE_DIR)], encoding="utf-8") + return int(line.split("\t")[0]) + + def _workspace_remove_latest(self): + # find oldest(lower number) directory + latest_dir = subprocess.check_output(self._adb_device_sub_cmd + [ + "shell", "find", str(ANDROID_HEXAGON_TEST_BASE_DIR), + "!", "-path", ".", "-type", "d", "|", "sort", "-n", "|", "head", "-1" + ], encoding="utf-8") + latest_dir = latest_dir.replace("\n", "").replace("\t", "") + + subprocess.check_call(self._adb_device_sub_cmd + ["shell", "rm", "-rf", latest_dir]) From b16c8b06dca8a83264dab0c7f7b4710f58f8ead9 Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Sat, 4 Dec 2021 02:22:01 -0800 Subject: [PATCH 04/15] Make libtvm_runtime consistent for Android --- cmake/modules/Hexagon.cmake | 2 +- python/tvm/contrib/hexagon/build.py | 4 +--- 2 files changed, 2 insertions(+), 4 deletions(-) diff --git a/cmake/modules/Hexagon.cmake b/cmake/modules/Hexagon.cmake index 6690c1a922ed..08f1fb0f35f6 100644 --- a/cmake/modules/Hexagon.cmake +++ b/cmake/modules/Hexagon.cmake @@ -252,7 +252,7 @@ if(USE_HEXAGON_RPC) ExternalProject_Add_Step(android_runtime_rpc copy_binary_runtime COMMAND ${CMAKE_COMMAND} -E copy_if_different ${BINARY_DIR}/libtvm_runtime.so - ${HEXAGON_RPC_OUTPUT}/libtvm_runtime_android.so + ${HEXAGON_RPC_OUTPUT}/libtvm_runtime.so DEPENDEES install ) ExternalProject_Add_Step(android_runtime_rpc copy_binary_rpc diff --git a/python/tvm/contrib/hexagon/build.py b/python/tvm/contrib/hexagon/build.py index 3d8602c02b1a..1306bbbb729c 100644 --- a/python/tvm/contrib/hexagon/build.py +++ b/python/tvm/contrib/hexagon/build.py @@ -29,7 +29,7 @@ from .session import Session -RPC_SERVER_FILES = ["tvm_rpc_android", "libtvm_runtime_android.so", "android_bash.sh"] +RPC_SERVER_FILES = ["tvm_rpc_android", "libtvm_runtime.so", "android_bash.sh"] HEXAGON_FILES = ["libhexagon_rpc_skel.so"] @@ -145,8 +145,6 @@ def android_run_rpc( for item in RPC_SERVER_FILES: src_path = get_hexagon_rpc_dir() / item destination = f"{self._workspace}/{item}" - if item == "libtvm_runtime_android.so": - destination = f"{self._workspace}/libtvm_runtime.so" subprocess.check_call(self._adb_device_sub_cmd + ["push", src_path, destination]) # enable root for adb From 076619f0d5a6cc91e3bc16882c2cd79743f40be9 Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Sat, 4 Dec 2021 02:25:17 -0800 Subject: [PATCH 05/15] Remove root access --- python/tvm/contrib/hexagon/build.py | 2 -- 1 file changed, 2 deletions(-) diff --git a/python/tvm/contrib/hexagon/build.py b/python/tvm/contrib/hexagon/build.py index 1306bbbb729c..34ff1004e3c6 100644 --- a/python/tvm/contrib/hexagon/build.py +++ b/python/tvm/contrib/hexagon/build.py @@ -147,8 +147,6 @@ def android_run_rpc( destination = f"{self._workspace}/{item}" subprocess.check_call(self._adb_device_sub_cmd + ["push", src_path, destination]) - # enable root for adb - subprocess.check_call(self._adb_device_sub_cmd + ["root"]) # Removed pre-defined forward/reverse rules subprocess.check_call(self._adb_device_sub_cmd + ["forward", "--remove-all"]) subprocess.check_call(self._adb_device_sub_cmd + ["reverse", "--remove-all"]) From 299a378d277bca48c1fc2ac2a5c9f79dd9052595 Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Sat, 4 Dec 2021 02:34:14 -0800 Subject: [PATCH 06/15] Fix some docstrings --- python/tvm/contrib/hexagon/build.py | 20 +++++++++---------- python/tvm/contrib/hexagon/session.py | 1 + .../contrib/test_hexagon/rpc/test_launcher.py | 2 +- 3 files changed, 12 insertions(+), 11 deletions(-) diff --git a/python/tvm/contrib/hexagon/build.py b/python/tvm/contrib/hexagon/build.py index 34ff1004e3c6..ae37f82aeadc 100644 --- a/python/tvm/contrib/hexagon/build.py +++ b/python/tvm/contrib/hexagon/build.py @@ -118,7 +118,7 @@ def android_run_rpc( # Upload RPC server and libraries subprocess.check_call(self._adb_device_sub_cmd + ["shell", "mkdir", "-p", self._workspace]) - # create bash script + # Create bash script android_bash_script_path = get_hexagon_rpc_dir() / "android_bash.sh" with open(get_hexagon_rpc_dir() / "android_bash.sh.template", "r") as src_f: if os.path.exists(android_bash_script_path): @@ -137,11 +137,11 @@ def android_run_rpc( line = line.replace("", str(rpc_server_port)) dest_f.write(line) - # make shell script executable + # Make shell script executable android_bash_stat = os.stat(android_bash_script_path) os.chmod(android_bash_script_path, android_bash_stat.st_mode | stat.S_IEXEC) - # push files + # Push files for item in RPC_SERVER_FILES: src_path = get_hexagon_rpc_dir() / item destination = f"{self._workspace}/{item}" @@ -226,16 +226,16 @@ def upload(self, host_path: Union[str, pathlib.Path], remote_filename: str): dst_remote_path = f"{self._workspace}/{remote_filename}" subprocess.check_call(self._adb_device_sub_cmd + ["push", src_path, dst_remote_path]) - def get_local_graph_executor(self, libmod, remote_libmod_filename: str): - """Create a GraphModule. + def get_graph_executor(self, libmod, remote_libmod_filename: str): + """Create a local GraphModule which consumes a remote libmod. Parameters ---------- libmod : tvm.runtime.Module - The module of the corresponding function + The module of the corresponding function. This library module is for remote hexagon runtime. remote_libmod_filename : str - Module filename on remote. Assumed this file lives under self._workspace path. + Module filename on remote. It is assumed this file lives under self._workspace path. Returns ------- @@ -250,11 +250,11 @@ def get_local_graph_executor(self, libmod, remote_libmod_filename: str): def close(self): """Close RPC server on Android""" - # kill process childs + # Kill process childs subprocess.Popen( self._adb_device_sub_cmd + ["shell", f"pkill -P `cat {self._workspace}/rpc_pid.txt`"] ) - # kill main process + # Kill main process subprocess.Popen( self._adb_device_sub_cmd + ["shell", f"kill `cat {self._workspace}/rpc_pid.txt`"] ) @@ -265,7 +265,7 @@ def _get_workspace_size(self) -> int: return int(line.split("\t")[0]) def _workspace_remove_latest(self): - # find oldest(lower number) directory + # Find oldest(lower number) directory latest_dir = subprocess.check_output(self._adb_device_sub_cmd + [ "shell", "find", str(ANDROID_HEXAGON_TEST_BASE_DIR), "!", "-path", ".", "-type", "d", "|", "sort", "-n", "|", "head", "-1" diff --git a/python/tvm/contrib/hexagon/session.py b/python/tvm/contrib/hexagon/session.py index 8cf449cbcc5a..25882a2526d9 100644 --- a/python/tvm/contrib/hexagon/session.py +++ b/python/tvm/contrib/hexagon/session.py @@ -44,6 +44,7 @@ def __init__( def __enter__(self): if self.device: + # Already initialized return self tracker = _rpc.connect_tracker(self._remote_kw["host"], self._remote_kw["port"]) diff --git a/tests/python/contrib/test_hexagon/rpc/test_launcher.py b/tests/python/contrib/test_hexagon/rpc/test_launcher.py index 1a1190b8f769..ff091bdc6e0a 100644 --- a/tests/python/contrib/test_hexagon/rpc/test_launcher.py +++ b/tests/python/contrib/test_hexagon/rpc/test_launcher.py @@ -185,7 +185,7 @@ def test_graph_executor(tvm_tracker_host, tvm_tracker_port, android_serial_numbe launcher.hexagon_session_setup(remote_kw) launcher.upload(dso_binary_path, dso_binary) - graph_mod = launcher.get_local_graph_executor(lowered, dso_binary) + graph_mod = launcher.get_graph_executor(lowered, dso_binary) weight_in = np.random.rand(5, 5, 3, 8).astype(dtype=dtype) data_in = np.random.rand(1, 64, 64, 3).astype(dtype=dtype) graph_mod.set_input(weight=weight_in) From 46562320069bb139e96bb86974b2a3cdbbc47c62 Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Mon, 6 Dec 2021 09:51:27 -0800 Subject: [PATCH 07/15] Make stack remote size as parameter --- python/tvm/contrib/hexagon/session.py | 3 +++ src/runtime/hexagon/rpc/android/session.cc | 7 ++++--- 2 files changed, 7 insertions(+), 3 deletions(-) diff --git a/python/tvm/contrib/hexagon/session.py b/python/tvm/contrib/hexagon/session.py index 25882a2526d9..3777a3192e5f 100644 --- a/python/tvm/contrib/hexagon/session.py +++ b/python/tvm/contrib/hexagon/session.py @@ -36,8 +36,10 @@ def __init__( self, remote_kw: dict, session_name: str = "hexagon-rpc", + remote_stack_size_bytes: int = 128 * 1024, ): self._session_name = session_name + self._remote_stack_size_bytes = remote_stack_size_bytes self._remote_kw = remote_kw self._rpc = None self.device = None @@ -56,6 +58,7 @@ def __enter__(self): session_constructor_args=[ "tvm.contrib.hexagon.create_hexagon_session", self._session_name, + self._remote_stack_size_bytes ], ) self.device = self._rpc.hexagon(0) diff --git a/src/runtime/hexagon/rpc/android/session.cc b/src/runtime/hexagon/rpc/android/session.cc index 886c1c7a22da..760b84ed0d8b 100644 --- a/src/runtime/hexagon/rpc/android/session.cc +++ b/src/runtime/hexagon/rpc/android/session.cc @@ -45,11 +45,11 @@ namespace hexagon { class HexagonTransportChannel : public RPCChannel { public: - explicit HexagonTransportChannel(const std::string& uri) { + explicit HexagonTransportChannel(const std::string& uri, int remote_stack_size_bytes) { if (_handle != AEE_EUNKNOWN) return; enable_unsigned_pd(true); - set_remote_stack_size(128 * 1024); + set_remote_stack_size(remote_stack_size_bytes); AEEResult rc = hexagon_rpc_open(uri.c_str(), &_handle); ICHECK(rc == AEE_SUCCESS) << "Hexagon RPC Open failed. URI: " << uri.c_str(); } @@ -106,8 +106,9 @@ class HexagonTransportChannel : public RPCChannel { TVM_REGISTER_GLOBAL("tvm.contrib.hexagon.create_hexagon_session") .set_body([](TVMArgs args, TVMRetValue* rv) { std::string session_name = args[0]; + int remote_stack_size_bytes = args[1]; HexagonTransportChannel* hexagon_channel = - new HexagonTransportChannel(hexagon_rpc_URI CDSP_DOMAIN); + new HexagonTransportChannel(hexagon_rpc_URI CDSP_DOMAIN, remote_stack_size_bytes); std::unique_ptr channel(hexagon_channel); auto ep = RPCEndpoint::Create(std::move(channel), session_name, "", NULL); auto sess = CreateClientSession(ep); From e57c7aafcc650a644e5d93d4ba7145725c6cd538 Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Mon, 6 Dec 2021 12:37:27 -0800 Subject: [PATCH 08/15] add documentation --- src/runtime/hexagon/rpc/hexagon/rpc_server.cc | 106 +++++++++++++----- src/runtime/hexagon/rpc/hexagon_rpc.idl | 2 +- 2 files changed, 77 insertions(+), 31 deletions(-) diff --git a/src/runtime/hexagon/rpc/hexagon/rpc_server.cc b/src/runtime/hexagon/rpc/hexagon/rpc_server.cc index b2f04cf4879b..6d5ee4dfc11c 100644 --- a/src/runtime/hexagon/rpc/hexagon/rpc_server.cc +++ b/src/runtime/hexagon/rpc/hexagon/rpc_server.cc @@ -48,18 +48,23 @@ namespace tvm { namespace runtime { namespace hexagon { +/*! + * \brief Hexagon IO Handler used in HexagonRPCServer(MinRPCServer). + * + * \param read_buffer The pointer to read buffer. + */ class HexagonIOHandler { public: - explicit HexagonIOHandler(uint8_t* read_buffer) : read_buffer_{read_buffer} {} + explicit HexagonIOHandler(uint8_t* read_buffer) : read_buffer_{read_buffer}, read_buffer_size_bytes_{0} {} void MessageStart(size_t message_size_bytes) {} ssize_t PosixWrite(const uint8_t* buf, size_t write_len_bytes) { - FARF(ALWAYS, "HexagonIOHandler PosixWrite called, write_len_bytes: %d", write_len_bytes); + HEXAGON_PRINT(ALWAYS, "HexagonIOHandler PosixWrite called, write_len_bytes: %d", write_len_bytes); size_t written_size = static_cast( write_buffer_.sputn(reinterpret_cast(buf), write_len_bytes)); if (written_size != write_len_bytes) { - FARF(ALWAYS, "HexagonIOHandler written_size failed"); + HEXAGON_PRINT(ALWAYS, "HexagonIOHandler written_size failed"); } return (ssize_t)written_size; } @@ -67,7 +72,7 @@ class HexagonIOHandler { void MessageDone() {} ssize_t PosixRead(uint8_t* buf, size_t read_len_bytes) { - FARF(ALWAYS, "HexagonIOHandler PosixRead called, %d, %d", read_len_bytes, + HEXAGON_PRINT(ALWAYS, "HexagonIOHandler PosixRead called, %d, %d", read_len_bytes, read_buffer_size_bytes_); uint32_t bytes_to_read = 0; @@ -81,24 +86,36 @@ class HexagonIOHandler { read_buffer_ += bytes_to_read; read_buffer_size_bytes_ -= bytes_to_read; if (bytes_to_read != read_len_bytes) { - FARF(ERROR, "Error bytes_to_read (%d) < read_len_bytes (%d).", bytes_to_read, read_len_bytes); + HEXAGON_PRINT(ERROR, "Error bytes_to_read (%d) < read_len_bytes (%d).", bytes_to_read, read_len_bytes); } return (ssize_t)bytes_to_read; } - void SetReadBuffer(const uint8_t* buf, size_t buf_size_bytes) { - FARF(ALWAYS, "HexagonIOHandler SetReadBuffer called: %d, prev read_buffer_size_bytes_: ", - buf_size_bytes, read_buffer_size_bytes_); - read_buffer_ = buf; - read_buffer_size_bytes_ = buf_size_bytes; + /*! + * \brief Set read buffer in IOHandler to data pointer. + * \param data The data pointer. + * \param data_size_bytes The size of data in bytes. + */ + void SetReadBuffer(const uint8_t* data, size_t data_size_bytes) { + HEXAGON_PRINT(ALWAYS, "HexagonIOHandler SetReadBuffer called: %d, prev read_buffer_size_bytes_: ", + data_size_bytes, read_buffer_size_bytes_); + read_buffer_ = data; + read_buffer_size_bytes_ = data_size_bytes; } - int64_t GetWriteBuffer(uint8_t* buf, size_t read_len_bytes) { - FARF(ALWAYS, "HexagonIOHandler GetWriteBuffer called, read_len_bytes: %d", read_len_bytes); - return write_buffer_.sgetn(reinterpret_cast(buf), read_len_bytes); + /*! + * \brief Get pointer to the buffer that a packet has been written to. + * \param buf The data pointer. + * \param read_size_bytes The size of read in bytes. + * + * \returns The size of data that is read in bytes. + */ + int64_t GetWriteBuffer(uint8_t* buf, size_t read_size_bytes) { + HEXAGON_PRINT(ALWAYS, "HexagonIOHandler GetWriteBuffer called, read_len_bytes: %d", read_size_bytes); + return write_buffer_.sgetn(reinterpret_cast(buf), read_size_bytes); } - void Close() { FARF(ALWAYS, "HexagonIOHandler Close called"); } + void Close() { HEXAGON_PRINT(ALWAYS, "HexagonIOHandler Close called"); } void Exit(int code) { exit(code); } @@ -113,14 +130,28 @@ class HexagonRPCServer { public: explicit HexagonRPCServer(uint8_t* receive_buffer) : io_{receive_buffer}, rpc_server_{&io_} {}; - int64_t Write(const uint8_t* data, size_t data_len_bytes) { - io_.SetReadBuffer(data, data_len_bytes); + /*! + * \brief Wrtie to IOHandler. + * \param data The data pointer + * \param data_size_bytes The data size in bytes. + * + * \returns The size of data written to IOHandler. + */ + int64_t Write(const uint8_t* data, size_t data_size_bytes) { + io_.SetReadBuffer(data, data_size_bytes); rpc_server_.ProcessOnePacket(); - return (int64_t)data_len_bytes; + return (int64_t)data_size_bytes; } - int64_t Read(uint8_t* buf, size_t read_len_bytes) { - return io_.GetWriteBuffer(buf, read_len_bytes); + /*! + * \brief Read from IOHandler. + * \param buf The buffer pointer + * \param read_size_bytes Read request size in bytes. + * + * \returns The size of data that is read in bytes. + */ + int64_t Read(uint8_t* buf, size_t read_size_bytes) { + return io_.GetWriteBuffer(buf, read_size_bytes); } private: @@ -156,7 +187,7 @@ void reset_device_api() { int __QAIC_HEADER(hexagon_rpc_open)(const char* uri, remote_handle64* handle) { *handle = static_cast(reinterpret_cast(malloc(1))); if (!*handle) { - FARF(ERROR, "%s: cannot allocate memory", __func__); + HEXAGON_PRINT(ERROR, "%s: cannot allocate memory", __func__); return AEE_ENOMEMORY; } reset_device_api(); @@ -165,41 +196,56 @@ int __QAIC_HEADER(hexagon_rpc_open)(const char* uri, remote_handle64* handle) { } int __QAIC_HEADER(hexagon_rpc_close)(remote_handle64 handle) { - FARF(ALWAYS, "%s", __func__); + HEXAGON_PRINT(ALWAYS, "%s", __func__); if (handle) { free(reinterpret_cast(static_cast(handle))); } return AEE_SUCCESS; } -// Send from Host to Hexagon over Android +/*! +* \brief Send data from Host to Hexagon over RPCSession. +* \param _handle The remote handle +* \param data The data sent to host. +* \param dataLen The size of the data. +* +* \returns The status. +*/ AEEResult __QAIC_HEADER(hexagon_rpc_send)(remote_handle64 _handle, const unsigned char* data, int dataLen) { if (g_hexagon_rpc_server == nullptr) { - FARF(ERROR, "RPC Server is not initialized."); + HEXAGON_PRINT(ERROR, "RPC Server is not initialized."); return AEE_EFAILED; } int64_t written_size = g_hexagon_rpc_server->Write(reinterpret_cast(data), static_cast(dataLen)); if (written_size != dataLen) { - FARF(ERROR, "RPC Server Write failed, written_size (%d) != dataLen (%d)", written_size, + HEXAGON_PRINT(ERROR, "RPC Server Write failed, written_size (%d) != dataLen (%d)", written_size, dataLen); return AEE_EFAILED; } return AEE_SUCCESS; } -// Receive from Hexagon and send to Host over Android. -AEEResult __QAIC_HEADER(hexagon_rpc_receive)(remote_handle64 _handle, unsigned char* data, - int dataLen, int64_t* buf_written_size) { +/*! +* \brief Receive data from Hexagon adn send to host over RPCSession. +* \param _handle The remote handle +* \param data The buffer for receiving data +* \param dataLen The size of the data that is requested to read in bytes. +* \param buf_written_size The size of the data that is actually read in bytes. +* +* \returns The status. +*/ +AEEResult __QAIC_HEADER(hexagon_rpc_receive)(remote_handle64 _handle, unsigned char* buf, + int bufLen, int64_t* buf_written_size) { int64_t read_size = - g_hexagon_rpc_server->Read(reinterpret_cast(data), static_cast(dataLen)); + g_hexagon_rpc_server->Read(reinterpret_cast(buf), static_cast(bufLen)); *buf_written_size = read_size; - if (read_size == dataLen) { + if (read_size == bufLen) { return AEE_SUCCESS; } else { - FARF(ALWAYS, "RPC Server Read failed, read_size (%d) != dataLen (%d)", read_size, dataLen); + HEXAGON_PRINT(ALWAYS, "RPC Server Read failed, read_size (%d) != dataLen (%d)", read_size, bufLen); return AEE_EFAILED; } } diff --git a/src/runtime/hexagon/rpc/hexagon_rpc.idl b/src/runtime/hexagon/rpc/hexagon_rpc.idl index 858ca501cf58..55b8d39bcb02 100644 --- a/src/runtime/hexagon/rpc/hexagon_rpc.idl +++ b/src/runtime/hexagon/rpc/hexagon_rpc.idl @@ -23,6 +23,6 @@ typedef sequence buffer; interface hexagon_rpc : remote_handle64 { - AEEResult send(in buffer buf); + AEEResult send(in buffer data); AEEResult receive(rout buffer buf, rout int64_t buf_written_size); }; From aa61967796f975156bcf61030a441a53c909ecf4 Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Mon, 6 Dec 2021 12:50:01 -0800 Subject: [PATCH 09/15] Refactor test conftest --- tests/python/contrib/test_hexagon/conftest.py | 99 +++++++++++++++++++ .../test_hexagon/proxy_rpc/test_matmul.py | 91 +---------------- .../contrib/test_hexagon/rpc/__init__.py | 0 .../contrib/test_hexagon/rpc/conftest.py | 54 +--------- .../contrib/test_hexagon/rpc/test_launcher.py | 3 +- 5 files changed, 103 insertions(+), 144 deletions(-) create mode 100644 tests/python/contrib/test_hexagon/rpc/__init__.py diff --git a/tests/python/contrib/test_hexagon/conftest.py b/tests/python/contrib/test_hexagon/conftest.py index b3bd00a08283..7028b0b30914 100644 --- a/tests/python/contrib/test_hexagon/conftest.py +++ b/tests/python/contrib/test_hexagon/conftest.py @@ -18,9 +18,108 @@ """ Hexagon testing fixtures used to deduce testing argument values from testing parameters """ +import os +import pytest + import tvm +from tvm import rpc +HEXAGON_TOOLCHAIN = "HEXAGON_TOOLCHAIN" +TVM_TRACKER_HOST = "TVM_TRACKER_HOST" +TVM_TRACKER_PORT = "TVM_TRACKER_PORT" +ANDROID_TRACKER_KEY = "ANDROID_TRACKER_KEY" +ANDROID_REMOTE_DIR = "ANDROID_REMOTE_DIR" @tvm.testing.fixture def shape_nhwc(batch, in_channel, in_size): return (batch, in_size, in_size, in_channel) + +def _compose(args, decs): + """Helper to apply multiple markers""" + if len(args) > 0: + f = args[0] + for d in reversed(decs): + f = d(f) + return f + return decs + +def requires_hexagon_toolchain(*args): + _requires_hexagon_toolchain = [ + pytest.mark.skipif( + os.environ.get("HEXAGON_TOOLCHAIN") == None, + reason="HEXAGON_TOOLCHAIN environment variable is required to run this test.", + ), + ] + + return _compose(args, _requires_hexagon_toolchain) + +@tvm.testing.fixture +def android_tracker_key(): + return os.environ["ANDROID_TRACKER_KEY"] + + +@tvm.testing.fixture +def tvm_tracker_host(): + return os.environ["TVM_TRACKER_HOST"] + + +@tvm.testing.fixture +def tvm_tracker_port(): + return int(os.environ["TVM_TRACKER_PORT"]) + + +@tvm.testing.fixture +def remote_path(): + dso_binary = "test_binary.so" + return os.path.join(os.environ["ANDROID_REMOTE_DIR"], dso_binary) + + +@tvm.testing.fixture +def rpc_sess(android_tracker_key, tvm_tracker_host, tvm_tracker_port): + from tvm import rpc + + tracker = rpc.connect_tracker(tvm_tracker_host, tvm_tracker_port) + remote = tracker.request(android_tracker_key, priority=0, session_timeout=600) + return remote + +def requires_rpc_tracker_and_android_key(*args): + """Mark a test as requiring an RPC tracker to exist in + the host environment to run.""" + _requires_rpc_tracker = [ + *tvm.testing.requires_rpc(), + pytest.mark.skipif( + os.environ.get(TVM_TRACKER_HOST) == None, + reason="Missing environment variable, TVM_TRACKER_HOST", + ), + pytest.mark.skipif( + os.environ.get(TVM_TRACKER_PORT) == None, + reason="Missing environment variable, TVM_TRACKER_PORT", + ), + pytest.mark.skipif( + os.environ.get(ANDROID_TRACKER_KEY) == None, + reason="Missing environment variable, ANDROID_TRACKER_KEY", + ), + pytest.mark.skipif( + os.environ.get(ANDROID_REMOTE_DIR) == None, + reason="Missing environment variable, ANDROID_REMOTE_DIR", + ), + ] + + return _compose(args, _requires_rpc_tracker) + +def requires_rpc_tracker(*args): + """Mark a test as requiring an RPC tracker to exist in + the host environment to run.""" + _requires_rpc_tracker = [ + *tvm.testing.requires_rpc(), + pytest.mark.skipif( + os.environ.get("TVM_TRACKER_HOST") == None, + reason="Missing environment variable, TVM_TRACKER_HOST", + ), + pytest.mark.skipif( + os.environ.get("TVM_TRACKER_PORT") == None, + reason="Missing environment variable, TVM_TRACKER_PORT", + ), + ] + + return _compose(args, _requires_rpc_tracker) diff --git a/tests/python/contrib/test_hexagon/proxy_rpc/test_matmul.py b/tests/python/contrib/test_hexagon/proxy_rpc/test_matmul.py index b83ce4480445..4e0483ce62ff 100644 --- a/tests/python/contrib/test_hexagon/proxy_rpc/test_matmul.py +++ b/tests/python/contrib/test_hexagon/proxy_rpc/test_matmul.py @@ -16,104 +16,17 @@ # under the License. import os -import sys import tvm import tvm.testing from tvm import te import tvm.contrib.hexagon.hexagon as hexagon from tvm.contrib import utils -from tvm import rpc import numpy as np -import pytest +from ..conftest import requires_hexagon_toolchain, requires_rpc_tracker_and_android_key -HEXAGON_TOOLCHAIN = "HEXAGON_TOOLCHAIN" -TVM_TRACKER_HOST = "TVM_TRACKER_HOST" -TVM_TRACKER_PORT = "TVM_TRACKER_PORT" -ANDROID_TRACKER_KEY = "ANDROID_TRACKER_KEY" -ANDROID_REMOTE_DIR = "ANDROID_REMOTE_DIR" - - -def _compose(args, decs): - """Helper to apply multiple markers""" - if len(args) > 0: - f = args[0] - for d in reversed(decs): - f = d(f) - return f - return decs - - -def requires_hexagon_toolchain(*args): - _requires_rpc_tracker = [ - *tvm.testing.requires_rpc(), - pytest.mark.skipif( - os.environ.get(HEXAGON_TOOLCHAIN) == None, - reason="HEXAGON_TOOLCHAIN environment variable is required to run Hexagon proxy rpc tests", - ), - ] - - return _compose(args, _requires_rpc_tracker) - - -def requires_rpc_tracker(*args): - """Mark a test as requiring an RPC tracker to exist in - the host environment to run.""" - _requires_rpc_tracker = [ - *tvm.testing.requires_rpc(), - pytest.mark.skipif( - os.environ.get(TVM_TRACKER_HOST) == None, - reason="Missing environment variable, TVM_TRACKER_HOST", - ), - pytest.mark.skipif( - os.environ.get(TVM_TRACKER_PORT) == None, - reason="Missing environment variable, TVM_TRACKER_PORT", - ), - pytest.mark.skipif( - os.environ.get(ANDROID_TRACKER_KEY) == None, - reason="Missing environment variable, ANDROID_TRACKER_KEY", - ), - pytest.mark.skipif( - os.environ.get(ANDROID_REMOTE_DIR) == None, - reason="Missing environment variable, ANDROID_REMOTE_DIR", - ), - ] - - return _compose(args, _requires_rpc_tracker) - - -@tvm.testing.fixture -def android_tracker_key(): - return os.environ["ANDROID_TRACKER_KEY"] - - -@tvm.testing.fixture -def tvm_tracker_host(): - return os.environ["TVM_TRACKER_HOST"] - - -@tvm.testing.fixture -def tvm_tracker_port(): - return int(os.environ["TVM_TRACKER_PORT"]) - - -@tvm.testing.fixture -def remote_path(): - dso_binary = "test_binary.so" - return os.path.join(os.environ["ANDROID_REMOTE_DIR"], dso_binary) - - -@tvm.testing.fixture -def rpc_sess(android_tracker_key, tvm_tracker_host, tvm_tracker_port): - from tvm import rpc - - tracker = rpc.connect_tracker(tvm_tracker_host, tvm_tracker_port) - remote = tracker.request(android_tracker_key, priority=0, session_timeout=600) - return remote - - -@requires_rpc_tracker +@requires_rpc_tracker_and_android_key @requires_hexagon_toolchain class TestMatMul: M = tvm.testing.parameter(32) diff --git a/tests/python/contrib/test_hexagon/rpc/__init__.py b/tests/python/contrib/test_hexagon/rpc/__init__.py new file mode 100644 index 000000000000..e69de29bb2d1 diff --git a/tests/python/contrib/test_hexagon/rpc/conftest.py b/tests/python/contrib/test_hexagon/rpc/conftest.py index 24edd0195504..08cf94d14132 100644 --- a/tests/python/contrib/test_hexagon/rpc/conftest.py +++ b/tests/python/contrib/test_hexagon/rpc/conftest.py @@ -14,11 +14,8 @@ # KIND, either express or implied. See the License for the # specific language governing permissions and limitations # under the License. -import pytest -import os - -import tvm.testing +import pytest def pytest_addoption(parser): parser.addoption( @@ -31,52 +28,3 @@ def pytest_addoption(parser): @pytest.fixture def android_serial_number(request): return request.config.getoption("--serial-number") - - -@tvm.testing.fixture -def tvm_tracker_host(): - return os.environ["TVM_TRACKER_HOST"] - - -@tvm.testing.fixture -def tvm_tracker_port(): - return int(os.environ["TVM_TRACKER_PORT"]) - - -def _compose(args, decs): - """Helper to apply multiple markers""" - if len(args) > 0: - f = args[0] - for d in reversed(decs): - f = d(f) - return f - return decs - - -def requires_rpc_tracker(*args): - """Mark a test as requiring an RPC tracker to exist in - the host environment to run.""" - _requires_rpc_tracker = [ - *tvm.testing.requires_rpc(), - pytest.mark.skipif( - os.environ.get("TVM_TRACKER_HOST") == None, - reason="Missing environment variable, TVM_TRACKER_HOST", - ), - pytest.mark.skipif( - os.environ.get("TVM_TRACKER_PORT") == None, - reason="Missing environment variable, TVM_TRACKER_PORT", - ), - ] - - return _compose(args, _requires_rpc_tracker) - - -def requires_hexagon_toolchain(*args): - _requires_hexagon_toolchain = [ - pytest.mark.skipif( - os.environ.get("HEXAGON_TOOLCHAIN") == None, - reason="HEXAGON_TOOLCHAIN environment variable is required to run this test.", - ), - ] - - return _compose(args, _requires_hexagon_toolchain) diff --git a/tests/python/contrib/test_hexagon/rpc/test_launcher.py b/tests/python/contrib/test_hexagon/rpc/test_launcher.py index ff091bdc6e0a..e86a133254f6 100644 --- a/tests/python/contrib/test_hexagon/rpc/test_launcher.py +++ b/tests/python/contrib/test_hexagon/rpc/test_launcher.py @@ -28,8 +28,7 @@ from tvm.contrib.hexagon.build import HexagonLauncher import tvm.contrib.hexagon.hexagon as hexagon -from conftest import requires_rpc_tracker, requires_hexagon_toolchain - +from ..conftest import requires_rpc_tracker, requires_hexagon_toolchain @requires_rpc_tracker @requires_hexagon_toolchain From 461df6a95de9a7ee0b6a21f16d87cb647149a54d Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Mon, 6 Dec 2021 12:58:55 -0800 Subject: [PATCH 10/15] clang format --- python/tvm/contrib/hexagon/build.py | 31 ++++++++-- python/tvm/contrib/hexagon/session.py | 2 +- src/runtime/hexagon/rpc/hexagon/rpc_server.cc | 60 ++++++++++--------- tests/python/contrib/test_hexagon/conftest.py | 6 ++ .../test_hexagon/proxy_rpc/test_matmul.py | 1 + .../contrib/test_hexagon/rpc/__init__.py | 18 ++++++ .../contrib/test_hexagon/rpc/conftest.py | 1 + .../contrib/test_hexagon/rpc/test_launcher.py | 1 + 8 files changed, 86 insertions(+), 34 deletions(-) diff --git a/python/tvm/contrib/hexagon/build.py b/python/tvm/contrib/hexagon/build.py index ae37f82aeadc..fe661e849199 100644 --- a/python/tvm/contrib/hexagon/build.py +++ b/python/tvm/contrib/hexagon/build.py @@ -106,7 +106,7 @@ def android_run_rpc( # Check size of base directory and cleanup if needed while self._get_workspace_size() > self._workspace_max_size_mb: self._workspace_remove_latest() - + if not workspace_dir: self._workspace = str( ANDROID_HEXAGON_TEST_BASE_DIR @@ -261,15 +261,34 @@ def close(self): def _get_workspace_size(self) -> int: """Get workspace base directory size in MB""" - line = subprocess.check_output(self._adb_device_sub_cmd + ["shell", "du", "-shm", str(ANDROID_HEXAGON_TEST_BASE_DIR)], encoding="utf-8") + line = subprocess.check_output( + self._adb_device_sub_cmd + ["shell", "du", "-shm", str(ANDROID_HEXAGON_TEST_BASE_DIR)], + encoding="utf-8", + ) return int(line.split("\t")[0]) def _workspace_remove_latest(self): # Find oldest(lower number) directory - latest_dir = subprocess.check_output(self._adb_device_sub_cmd + [ - "shell", "find", str(ANDROID_HEXAGON_TEST_BASE_DIR), - "!", "-path", ".", "-type", "d", "|", "sort", "-n", "|", "head", "-1" - ], encoding="utf-8") + latest_dir = subprocess.check_output( + self._adb_device_sub_cmd + + [ + "shell", + "find", + str(ANDROID_HEXAGON_TEST_BASE_DIR), + "!", + "-path", + ".", + "-type", + "d", + "|", + "sort", + "-n", + "|", + "head", + "-1", + ], + encoding="utf-8", + ) latest_dir = latest_dir.replace("\n", "").replace("\t", "") subprocess.check_call(self._adb_device_sub_cmd + ["shell", "rm", "-rf", latest_dir]) diff --git a/python/tvm/contrib/hexagon/session.py b/python/tvm/contrib/hexagon/session.py index 3777a3192e5f..c413c60ce5c5 100644 --- a/python/tvm/contrib/hexagon/session.py +++ b/python/tvm/contrib/hexagon/session.py @@ -58,7 +58,7 @@ def __enter__(self): session_constructor_args=[ "tvm.contrib.hexagon.create_hexagon_session", self._session_name, - self._remote_stack_size_bytes + self._remote_stack_size_bytes, ], ) self.device = self._rpc.hexagon(0) diff --git a/src/runtime/hexagon/rpc/hexagon/rpc_server.cc b/src/runtime/hexagon/rpc/hexagon/rpc_server.cc index 6d5ee4dfc11c..1f9c7c7ce82b 100644 --- a/src/runtime/hexagon/rpc/hexagon/rpc_server.cc +++ b/src/runtime/hexagon/rpc/hexagon/rpc_server.cc @@ -55,12 +55,14 @@ namespace hexagon { */ class HexagonIOHandler { public: - explicit HexagonIOHandler(uint8_t* read_buffer) : read_buffer_{read_buffer}, read_buffer_size_bytes_{0} {} + explicit HexagonIOHandler(uint8_t* read_buffer) + : read_buffer_{read_buffer}, read_buffer_size_bytes_{0} {} void MessageStart(size_t message_size_bytes) {} ssize_t PosixWrite(const uint8_t* buf, size_t write_len_bytes) { - HEXAGON_PRINT(ALWAYS, "HexagonIOHandler PosixWrite called, write_len_bytes: %d", write_len_bytes); + HEXAGON_PRINT(ALWAYS, "HexagonIOHandler PosixWrite called, write_len_bytes: %d", + write_len_bytes); size_t written_size = static_cast( write_buffer_.sputn(reinterpret_cast(buf), write_len_bytes)); if (written_size != write_len_bytes) { @@ -73,7 +75,7 @@ class HexagonIOHandler { ssize_t PosixRead(uint8_t* buf, size_t read_len_bytes) { HEXAGON_PRINT(ALWAYS, "HexagonIOHandler PosixRead called, %d, %d", read_len_bytes, - read_buffer_size_bytes_); + read_buffer_size_bytes_); uint32_t bytes_to_read = 0; if ((read_buffer_size_bytes_ - read_len_bytes) < 0) { @@ -86,7 +88,8 @@ class HexagonIOHandler { read_buffer_ += bytes_to_read; read_buffer_size_bytes_ -= bytes_to_read; if (bytes_to_read != read_len_bytes) { - HEXAGON_PRINT(ERROR, "Error bytes_to_read (%d) < read_len_bytes (%d).", bytes_to_read, read_len_bytes); + HEXAGON_PRINT(ERROR, "Error bytes_to_read (%d) < read_len_bytes (%d).", bytes_to_read, + read_len_bytes); } return (ssize_t)bytes_to_read; } @@ -97,8 +100,9 @@ class HexagonIOHandler { * \param data_size_bytes The size of data in bytes. */ void SetReadBuffer(const uint8_t* data, size_t data_size_bytes) { - HEXAGON_PRINT(ALWAYS, "HexagonIOHandler SetReadBuffer called: %d, prev read_buffer_size_bytes_: ", - data_size_bytes, read_buffer_size_bytes_); + HEXAGON_PRINT(ALWAYS, + "HexagonIOHandler SetReadBuffer called: %d, prev read_buffer_size_bytes_: ", + data_size_bytes, read_buffer_size_bytes_); read_buffer_ = data; read_buffer_size_bytes_ = data_size_bytes; } @@ -107,11 +111,12 @@ class HexagonIOHandler { * \brief Get pointer to the buffer that a packet has been written to. * \param buf The data pointer. * \param read_size_bytes The size of read in bytes. - * + * * \returns The size of data that is read in bytes. */ int64_t GetWriteBuffer(uint8_t* buf, size_t read_size_bytes) { - HEXAGON_PRINT(ALWAYS, "HexagonIOHandler GetWriteBuffer called, read_len_bytes: %d", read_size_bytes); + HEXAGON_PRINT(ALWAYS, "HexagonIOHandler GetWriteBuffer called, read_len_bytes: %d", + read_size_bytes); return write_buffer_.sgetn(reinterpret_cast(buf), read_size_bytes); } @@ -134,7 +139,7 @@ class HexagonRPCServer { * \brief Wrtie to IOHandler. * \param data The data pointer * \param data_size_bytes The data size in bytes. - * + * * \returns The size of data written to IOHandler. */ int64_t Write(const uint8_t* data, size_t data_size_bytes) { @@ -147,7 +152,7 @@ class HexagonRPCServer { * \brief Read from IOHandler. * \param buf The buffer pointer * \param read_size_bytes Read request size in bytes. - * + * * \returns The size of data that is read in bytes. */ int64_t Read(uint8_t* buf, size_t read_size_bytes) { @@ -204,13 +209,13 @@ int __QAIC_HEADER(hexagon_rpc_close)(remote_handle64 handle) { } /*! -* \brief Send data from Host to Hexagon over RPCSession. -* \param _handle The remote handle -* \param data The data sent to host. -* \param dataLen The size of the data. -* -* \returns The status. -*/ + * \brief Send data from Host to Hexagon over RPCSession. + * \param _handle The remote handle + * \param data The data sent to host. + * \param dataLen The size of the data. + * + * \returns The status. + */ AEEResult __QAIC_HEADER(hexagon_rpc_send)(remote_handle64 _handle, const unsigned char* data, int dataLen) { if (g_hexagon_rpc_server == nullptr) { @@ -222,21 +227,21 @@ AEEResult __QAIC_HEADER(hexagon_rpc_send)(remote_handle64 _handle, const unsigne static_cast(dataLen)); if (written_size != dataLen) { HEXAGON_PRINT(ERROR, "RPC Server Write failed, written_size (%d) != dataLen (%d)", written_size, - dataLen); + dataLen); return AEE_EFAILED; } return AEE_SUCCESS; } /*! -* \brief Receive data from Hexagon adn send to host over RPCSession. -* \param _handle The remote handle -* \param data The buffer for receiving data -* \param dataLen The size of the data that is requested to read in bytes. -* \param buf_written_size The size of the data that is actually read in bytes. -* -* \returns The status. -*/ + * \brief Receive data from Hexagon adn send to host over RPCSession. + * \param _handle The remote handle + * \param data The buffer for receiving data + * \param dataLen The size of the data that is requested to read in bytes. + * \param buf_written_size The size of the data that is actually read in bytes. + * + * \returns The status. + */ AEEResult __QAIC_HEADER(hexagon_rpc_receive)(remote_handle64 _handle, unsigned char* buf, int bufLen, int64_t* buf_written_size) { int64_t read_size = @@ -245,7 +250,8 @@ AEEResult __QAIC_HEADER(hexagon_rpc_receive)(remote_handle64 _handle, unsigned c if (read_size == bufLen) { return AEE_SUCCESS; } else { - HEXAGON_PRINT(ALWAYS, "RPC Server Read failed, read_size (%d) != dataLen (%d)", read_size, bufLen); + HEXAGON_PRINT(ALWAYS, "RPC Server Read failed, read_size (%d) != dataLen (%d)", read_size, + bufLen); return AEE_EFAILED; } } diff --git a/tests/python/contrib/test_hexagon/conftest.py b/tests/python/contrib/test_hexagon/conftest.py index 7028b0b30914..bd2ae7c08473 100644 --- a/tests/python/contrib/test_hexagon/conftest.py +++ b/tests/python/contrib/test_hexagon/conftest.py @@ -30,10 +30,12 @@ ANDROID_TRACKER_KEY = "ANDROID_TRACKER_KEY" ANDROID_REMOTE_DIR = "ANDROID_REMOTE_DIR" + @tvm.testing.fixture def shape_nhwc(batch, in_channel, in_size): return (batch, in_size, in_size, in_channel) + def _compose(args, decs): """Helper to apply multiple markers""" if len(args) > 0: @@ -43,6 +45,7 @@ def _compose(args, decs): return f return decs + def requires_hexagon_toolchain(*args): _requires_hexagon_toolchain = [ pytest.mark.skipif( @@ -53,6 +56,7 @@ def requires_hexagon_toolchain(*args): return _compose(args, _requires_hexagon_toolchain) + @tvm.testing.fixture def android_tracker_key(): return os.environ["ANDROID_TRACKER_KEY"] @@ -82,6 +86,7 @@ def rpc_sess(android_tracker_key, tvm_tracker_host, tvm_tracker_port): remote = tracker.request(android_tracker_key, priority=0, session_timeout=600) return remote + def requires_rpc_tracker_and_android_key(*args): """Mark a test as requiring an RPC tracker to exist in the host environment to run.""" @@ -107,6 +112,7 @@ def requires_rpc_tracker_and_android_key(*args): return _compose(args, _requires_rpc_tracker) + def requires_rpc_tracker(*args): """Mark a test as requiring an RPC tracker to exist in the host environment to run.""" diff --git a/tests/python/contrib/test_hexagon/proxy_rpc/test_matmul.py b/tests/python/contrib/test_hexagon/proxy_rpc/test_matmul.py index 4e0483ce62ff..839fdc9bc29d 100644 --- a/tests/python/contrib/test_hexagon/proxy_rpc/test_matmul.py +++ b/tests/python/contrib/test_hexagon/proxy_rpc/test_matmul.py @@ -26,6 +26,7 @@ from ..conftest import requires_hexagon_toolchain, requires_rpc_tracker_and_android_key + @requires_rpc_tracker_and_android_key @requires_hexagon_toolchain class TestMatMul: diff --git a/tests/python/contrib/test_hexagon/rpc/__init__.py b/tests/python/contrib/test_hexagon/rpc/__init__.py index e69de29bb2d1..92e96bf52e8b 100644 --- a/tests/python/contrib/test_hexagon/rpc/__init__.py +++ b/tests/python/contrib/test_hexagon/rpc/__init__.py @@ -0,0 +1,18 @@ +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you under the Apache License, Version 2.0 (the +# "License"); you may not use this file except in compliance +# with the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, +# software distributed under the License is distributed on an +# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +# KIND, either express or implied. See the License for the +# specific language governing permissions and limitations +# under the License. + +""" Testing infrastructure for Hexagon RPC""" diff --git a/tests/python/contrib/test_hexagon/rpc/conftest.py b/tests/python/contrib/test_hexagon/rpc/conftest.py index 08cf94d14132..50c199a2f56f 100644 --- a/tests/python/contrib/test_hexagon/rpc/conftest.py +++ b/tests/python/contrib/test_hexagon/rpc/conftest.py @@ -17,6 +17,7 @@ import pytest + def pytest_addoption(parser): parser.addoption( "--serial-number", diff --git a/tests/python/contrib/test_hexagon/rpc/test_launcher.py b/tests/python/contrib/test_hexagon/rpc/test_launcher.py index e86a133254f6..d705541f2b8e 100644 --- a/tests/python/contrib/test_hexagon/rpc/test_launcher.py +++ b/tests/python/contrib/test_hexagon/rpc/test_launcher.py @@ -30,6 +30,7 @@ from ..conftest import requires_rpc_tracker, requires_hexagon_toolchain + @requires_rpc_tracker @requires_hexagon_toolchain def test_add(tvm_tracker_host, tvm_tracker_port, android_serial_number): From 28db608e66b16ffb8159113d2edd4408a90caa89 Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Mon, 6 Dec 2021 15:56:38 -0800 Subject: [PATCH 11/15] Decoupled USE_HEXAGON_RPC --- cmake/modules/Hexagon.cmake | 10 +++++++--- 1 file changed, 7 insertions(+), 3 deletions(-) diff --git a/cmake/modules/Hexagon.cmake b/cmake/modules/Hexagon.cmake index 08f1fb0f35f6..b9623f3fcec8 100644 --- a/cmake/modules/Hexagon.cmake +++ b/cmake/modules/Hexagon.cmake @@ -293,6 +293,8 @@ if(USE_HEXAGON_RPC) # copy android_bash template file configure_file("${CMAKE_SOURCE_DIR}/src/runtime/hexagon/rpc/android_bash.sh.template" ${HEXAGON_RPC_OUTPUT} COPYONLY) + + set_directory_properties(PROPERTIES ADDITIONAL_MAKE_CLEAN_FILES "${HEXAGON_RPC_OUTPUT}") endif() if (USE_HEXAGON_SDK AND BUILD_FOR_ANDROID) @@ -317,7 +319,7 @@ if(USE_HEXAGON_DEVICE STREQUAL "${PICK_SIM}") "-DHEXAGON_ARCH=${USE_HEXAGON_ARCH}" INSTALL_COMMAND "true" ) -elseif((USE_HEXAGON_DEVICE STREQUAL "${PICK_HW}") OR (USE_HEXAGON_RPC AND BUILD_FOR_ANDROID)) +elseif(USE_HEXAGON_DEVICE STREQUAL "${PICK_HW}") find_hexagon_sdk_root("${USE_HEXAGON_SDK}" "${USE_HEXAGON_ARCH}") find_hexagon_toolchain() @@ -344,13 +346,15 @@ if (USE_HEXAGON_DEVICE STREQUAL "${PICK_NONE}") file(GLOB RUNTIME_HEXAGON_SRCS src/runtime/hexagon/hexagon/*.cc) elseif(BUILD_FOR_ANDROID AND HEXAGON_SDK_PATH_DEFINED) list(APPEND RUNTIME_HEXAGON_SRCS src/runtime/hexagon/proxy_rpc/device_api.cc) - elseif(USE_HEXAGON_RPC) - file(GLOB RUNTIME_HEXAGON_SRCS src/runtime/hexagon/host/*.cc) endif() else() file(GLOB RUNTIME_HEXAGON_SRCS src/runtime/hexagon/android/*.cc) endif() +if(USE_HEXAGON_RPC) + file(GLOB RUNTIME_HEXAGON_SRCS src/runtime/hexagon/host/*.cc) +endif() + if(USE_HEXAGON_SDK AND BUILD_FOR_ANDROID) find_hexagon_sdk_root("${USE_HEXAGON_SDK}" "${USE_HEXAGON_ARCH}") include_directories(SYSTEM ${HEXAGON_SDK_INCLUDES} ${HEXAGON_REMOTE_ROOT}) From f87f51d5f61684b72bf06d79bdc2f40f20bd7191 Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Mon, 6 Dec 2021 16:05:39 -0800 Subject: [PATCH 12/15] fix creation of test base directory on android --- python/tvm/contrib/hexagon/build.py | 3 +++ 1 file changed, 3 insertions(+) diff --git a/python/tvm/contrib/hexagon/build.py b/python/tvm/contrib/hexagon/build.py index fe661e849199..816f36662b6f 100644 --- a/python/tvm/contrib/hexagon/build.py +++ b/python/tvm/contrib/hexagon/build.py @@ -103,6 +103,9 @@ def android_run_rpc( rpc_tracker_port : int RPC tracker port on host """ + # Create test base directory + subprocess.check_call(self._adb_device_sub_cmd + ["shell", "mkdir", "-p", ANDROID_HEXAGON_TEST_BASE_DIR]) + # Check size of base directory and cleanup if needed while self._get_workspace_size() > self._workspace_max_size_mb: self._workspace_remove_latest() From 31fda2bd91f0f73b99bde0b85f32180b9404c58a Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Tue, 7 Dec 2021 15:09:09 -0800 Subject: [PATCH 13/15] Address global variable --- src/runtime/hexagon/rpc/hexagon/rpc_server.cc | 69 ++++++++++--------- 1 file changed, 36 insertions(+), 33 deletions(-) diff --git a/src/runtime/hexagon/rpc/hexagon/rpc_server.cc b/src/runtime/hexagon/rpc/hexagon/rpc_server.cc index 1f9c7c7ce82b..bd9de21e20cf 100644 --- a/src/runtime/hexagon/rpc/hexagon/rpc_server.cc +++ b/src/runtime/hexagon/rpc/hexagon/rpc_server.cc @@ -52,11 +52,12 @@ namespace hexagon { * \brief Hexagon IO Handler used in HexagonRPCServer(MinRPCServer). * * \param read_buffer The pointer to read buffer. + * \param read_buffer_size_bytes The read buffer size in bytes. */ class HexagonIOHandler { public: - explicit HexagonIOHandler(uint8_t* read_buffer) - : read_buffer_{read_buffer}, read_buffer_size_bytes_{0} {} + explicit HexagonIOHandler(uint8_t* read_buffer, size_t read_buffer_size_bytes) + : read_buffer_{read_buffer}, read_buffer_size_bytes_{read_buffer_size_bytes}, read_buffer_index_{0} {} void MessageStart(size_t message_size_bytes) {} @@ -75,18 +76,18 @@ class HexagonIOHandler { ssize_t PosixRead(uint8_t* buf, size_t read_len_bytes) { HEXAGON_PRINT(ALWAYS, "HexagonIOHandler PosixRead called, %d, %d", read_len_bytes, - read_buffer_size_bytes_); + read_buffer_index_); uint32_t bytes_to_read = 0; - if ((read_buffer_size_bytes_ - read_len_bytes) < 0) { - bytes_to_read = read_buffer_size_bytes_; + if ((read_buffer_index_ - read_len_bytes) < 0) { + bytes_to_read = read_buffer_index_; } else { bytes_to_read = read_len_bytes; } std::memcpy(buf, read_buffer_, bytes_to_read); read_buffer_ += bytes_to_read; - read_buffer_size_bytes_ -= bytes_to_read; + read_buffer_index_ -= bytes_to_read; if (bytes_to_read != read_len_bytes) { HEXAGON_PRINT(ERROR, "Error bytes_to_read (%d) < read_len_bytes (%d).", bytes_to_read, read_len_bytes); @@ -98,13 +99,19 @@ class HexagonIOHandler { * \brief Set read buffer in IOHandler to data pointer. * \param data The data pointer. * \param data_size_bytes The size of data in bytes. + * + * \return The status */ - void SetReadBuffer(const uint8_t* data, size_t data_size_bytes) { + AEEResult SetReadBuffer(const uint8_t* data, size_t data_size_bytes) { HEXAGON_PRINT(ALWAYS, - "HexagonIOHandler SetReadBuffer called: %d, prev read_buffer_size_bytes_: ", - data_size_bytes, read_buffer_size_bytes_); + "HexagonIOHandler SetReadBuffer called: %d, prev read_buffer_index_: ", + data_size_bytes, read_buffer_index_); + if (data_size_bytes > read_buffer_size_bytes_) { + return AEE_EFAILED; + } read_buffer_ = data; - read_buffer_size_bytes_ = data_size_bytes; + read_buffer_index_ = data_size_bytes; + return AEE_SUCCESS; } /*! @@ -112,7 +119,7 @@ class HexagonIOHandler { * \param buf The data pointer. * \param read_size_bytes The size of read in bytes. * - * \returns The size of data that is read in bytes. + * \return The size of data that is read in bytes. */ int64_t GetWriteBuffer(uint8_t* buf, size_t read_size_bytes) { HEXAGON_PRINT(ALWAYS, "HexagonIOHandler GetWriteBuffer called, read_len_bytes: %d", @@ -126,24 +133,27 @@ class HexagonIOHandler { private: const uint8_t* read_buffer_; - uint32_t read_buffer_size_bytes_; + uint32_t read_buffer_index_; + size_t read_buffer_size_bytes_; std::stringbuf write_buffer_; }; class HexagonRPCServer { public: - explicit HexagonRPCServer(uint8_t* receive_buffer) : io_{receive_buffer}, rpc_server_{&io_} {}; + explicit HexagonRPCServer(uint8_t* receive_buffer, size_t receive_buffer_size_bytes) : io_{receive_buffer, receive_buffer_size_bytes}, rpc_server_{&io_} {}; /*! * \brief Wrtie to IOHandler. * \param data The data pointer * \param data_size_bytes The data size in bytes. * - * \returns The size of data written to IOHandler. + * \return The size of data written to IOHandler. */ int64_t Write(const uint8_t* data, size_t data_size_bytes) { - io_.SetReadBuffer(data, data_size_bytes); + if (io_.SetReadBuffer(data, data_size_bytes) != AEE_SUCCESS) { + return -1; + } rpc_server_.ProcessOnePacket(); return (int64_t)data_size_bytes; } @@ -153,7 +163,7 @@ class HexagonRPCServer { * \param buf The buffer pointer * \param read_size_bytes Read request size in bytes. * - * \returns The size of data that is read in bytes. + * \return The size of data that is read in bytes. */ int64_t Read(uint8_t* buf, size_t read_size_bytes) { return io_.GetWriteBuffer(buf, read_size_bytes); @@ -168,13 +178,11 @@ class HexagonRPCServer { } // namespace runtime } // namespace tvm -static tvm::runtime::hexagon::HexagonRPCServer* g_hexagon_rpc_server = nullptr; - -static AEEResult hexagon_rpc_server_init() { - uint8_t* receive_buffer = new uint8_t[TVM_HEXAGON_RPC_BUFF_SIZE_BYTES]; - tvm::runtime::hexagon::HexagonRPCServer* rpc_server = - new tvm::runtime::hexagon::HexagonRPCServer(receive_buffer); - g_hexagon_rpc_server = rpc_server; +namespace { + tvm::runtime::hexagon::HexagonRPCServer* get_hexagon_rpc_server() { + static tvm::runtime::hexagon::HexagonRPCServer g_hexagon_rpc_server(new uint8_t[TVM_HEXAGON_RPC_BUFF_SIZE_BYTES], TVM_HEXAGON_RPC_BUFF_SIZE_BYTES); + return &g_hexagon_rpc_server; + } } const tvm::runtime::PackedFunc get_runtime_func(const std::string& name) { @@ -196,7 +204,7 @@ int __QAIC_HEADER(hexagon_rpc_open)(const char* uri, remote_handle64* handle) { return AEE_ENOMEMORY; } reset_device_api(); - hexagon_rpc_server_init(); + get_hexagon_rpc_server(); return AEE_SUCCESS; } @@ -214,16 +222,11 @@ int __QAIC_HEADER(hexagon_rpc_close)(remote_handle64 handle) { * \param data The data sent to host. * \param dataLen The size of the data. * - * \returns The status. + * \return The status. */ AEEResult __QAIC_HEADER(hexagon_rpc_send)(remote_handle64 _handle, const unsigned char* data, int dataLen) { - if (g_hexagon_rpc_server == nullptr) { - HEXAGON_PRINT(ERROR, "RPC Server is not initialized."); - return AEE_EFAILED; - } - - int64_t written_size = g_hexagon_rpc_server->Write(reinterpret_cast(data), + int64_t written_size = get_hexagon_rpc_server()->Write(reinterpret_cast(data), static_cast(dataLen)); if (written_size != dataLen) { HEXAGON_PRINT(ERROR, "RPC Server Write failed, written_size (%d) != dataLen (%d)", written_size, @@ -240,12 +243,12 @@ AEEResult __QAIC_HEADER(hexagon_rpc_send)(remote_handle64 _handle, const unsigne * \param dataLen The size of the data that is requested to read in bytes. * \param buf_written_size The size of the data that is actually read in bytes. * - * \returns The status. + * \return The status. */ AEEResult __QAIC_HEADER(hexagon_rpc_receive)(remote_handle64 _handle, unsigned char* buf, int bufLen, int64_t* buf_written_size) { int64_t read_size = - g_hexagon_rpc_server->Read(reinterpret_cast(buf), static_cast(bufLen)); + get_hexagon_rpc_server()->Read(reinterpret_cast(buf), static_cast(bufLen)); *buf_written_size = read_size; if (read_size == bufLen) { return AEE_SUCCESS; From a7ee1a9578551d02d86f93aa129c7b781d146b3a Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Tue, 7 Dec 2021 15:17:06 -0800 Subject: [PATCH 14/15] Fix format and Cleanup cmake --- cmake/modules/Hexagon.cmake | 14 +++-------- python/tvm/contrib/hexagon/__init__.py | 1 + python/tvm/contrib/hexagon/build.py | 7 ++++-- src/runtime/hexagon/rpc/hexagon/rpc_server.cc | 25 +++++++++++-------- 4 files changed, 23 insertions(+), 24 deletions(-) diff --git a/cmake/modules/Hexagon.cmake b/cmake/modules/Hexagon.cmake index b9623f3fcec8..7d33b6b58d53 100644 --- a/cmake/modules/Hexagon.cmake +++ b/cmake/modules/Hexagon.cmake @@ -297,13 +297,6 @@ if(USE_HEXAGON_RPC) set_directory_properties(PROPERTIES ADDITIONAL_MAKE_CLEAN_FILES "${HEXAGON_RPC_OUTPUT}") endif() -if (USE_HEXAGON_SDK AND BUILD_FOR_ANDROID) - find_hexagon_sdk_root("${USE_HEXAGON_SDK}" "${USE_HEXAGON_ARCH}") - include_directories(SYSTEM ${HEXAGON_SDK_INCLUDES} ${HEXAGON_REMOTE_ROOT}) - link_directories(${HEXAGON_REMOTE_ROOT}) - list(APPEND TVM_RUNTIME_LINKER_LIBS cdsprpc) -endif() - if(USE_HEXAGON_DEVICE STREQUAL "${PICK_SIM}") find_hexagon_toolchain() message(STATUS "Hexagon toolchain: ${HEXAGON_TOOLCHAIN}") @@ -322,10 +315,7 @@ if(USE_HEXAGON_DEVICE STREQUAL "${PICK_SIM}") elseif(USE_HEXAGON_DEVICE STREQUAL "${PICK_HW}") find_hexagon_sdk_root("${USE_HEXAGON_SDK}" "${USE_HEXAGON_ARCH}") find_hexagon_toolchain() - - if(USE_HEXAGON_DEVICE STREQUAL "${PICK_HW}") - file(GLOB RUNTIME_HEXAGON_DEVICE_SRCS src/runtime/hexagon/android/target/*.cc) - endif() + file(GLOB RUNTIME_HEXAGON_DEVICE_SRCS src/runtime/hexagon/android/target/*.cc) include_directories(SYSTEM ${HEXAGON_SDK_INCLUDES} @@ -346,6 +336,8 @@ if (USE_HEXAGON_DEVICE STREQUAL "${PICK_NONE}") file(GLOB RUNTIME_HEXAGON_SRCS src/runtime/hexagon/hexagon/*.cc) elseif(BUILD_FOR_ANDROID AND HEXAGON_SDK_PATH_DEFINED) list(APPEND RUNTIME_HEXAGON_SRCS src/runtime/hexagon/proxy_rpc/device_api.cc) + else() + file(GLOB RUNTIME_HEXAGON_SRCS src/runtime/hexagon/host/*.cc) endif() else() file(GLOB RUNTIME_HEXAGON_SRCS src/runtime/hexagon/android/*.cc) diff --git a/python/tvm/contrib/hexagon/__init__.py b/python/tvm/contrib/hexagon/__init__.py index 13a83393a912..e728216853e5 100644 --- a/python/tvm/contrib/hexagon/__init__.py +++ b/python/tvm/contrib/hexagon/__init__.py @@ -14,3 +14,4 @@ # KIND, either express or implied. See the License for the # specific language governing permissions and limitations # under the License. +"""Hexagon APIs.""" diff --git a/python/tvm/contrib/hexagon/build.py b/python/tvm/contrib/hexagon/build.py index 816f36662b6f..ef081f2d79e3 100644 --- a/python/tvm/contrib/hexagon/build.py +++ b/python/tvm/contrib/hexagon/build.py @@ -104,7 +104,9 @@ def android_run_rpc( RPC tracker port on host """ # Create test base directory - subprocess.check_call(self._adb_device_sub_cmd + ["shell", "mkdir", "-p", ANDROID_HEXAGON_TEST_BASE_DIR]) + subprocess.check_call( + self._adb_device_sub_cmd + ["shell", "mkdir", "-p", ANDROID_HEXAGON_TEST_BASE_DIR] + ) # Check size of base directory and cleanup if needed while self._get_workspace_size() > self._workspace_max_size_mb: @@ -235,7 +237,8 @@ def get_graph_executor(self, libmod, remote_libmod_filename: str): Parameters ---------- libmod : tvm.runtime.Module - The module of the corresponding function. This library module is for remote hexagon runtime. + The module of the corresponding function. + This library module is for remote hexagon runtime. remote_libmod_filename : str Module filename on remote. It is assumed this file lives under self._workspace path. diff --git a/src/runtime/hexagon/rpc/hexagon/rpc_server.cc b/src/runtime/hexagon/rpc/hexagon/rpc_server.cc index bd9de21e20cf..bf4b0e3f5d85 100644 --- a/src/runtime/hexagon/rpc/hexagon/rpc_server.cc +++ b/src/runtime/hexagon/rpc/hexagon/rpc_server.cc @@ -40,7 +40,7 @@ extern "C" { #include "../../hexagon/hexagon_common.h" #include "hexagon_rpc.h" -#define TVM_HEXAGON_RPC_BUFF_SIZE_BYTES 1024 * 1024 +#define TVM_HEXAGON_RPC_BUFF_SIZE_BYTES 2 * 1024 * 1024 #define TVM_LOG_CUSTOMIZE 1 @@ -57,7 +57,9 @@ namespace hexagon { class HexagonIOHandler { public: explicit HexagonIOHandler(uint8_t* read_buffer, size_t read_buffer_size_bytes) - : read_buffer_{read_buffer}, read_buffer_size_bytes_{read_buffer_size_bytes}, read_buffer_index_{0} {} + : read_buffer_{read_buffer}, + read_buffer_size_bytes_{read_buffer_size_bytes}, + read_buffer_index_{0} {} void MessageStart(size_t message_size_bytes) {} @@ -99,12 +101,11 @@ class HexagonIOHandler { * \brief Set read buffer in IOHandler to data pointer. * \param data The data pointer. * \param data_size_bytes The size of data in bytes. - * + * * \return The status */ AEEResult SetReadBuffer(const uint8_t* data, size_t data_size_bytes) { - HEXAGON_PRINT(ALWAYS, - "HexagonIOHandler SetReadBuffer called: %d, prev read_buffer_index_: ", + HEXAGON_PRINT(ALWAYS, "HexagonIOHandler SetReadBuffer called: %d, prev read_buffer_index_: ", data_size_bytes, read_buffer_index_); if (data_size_bytes > read_buffer_size_bytes_) { return AEE_EFAILED; @@ -141,7 +142,8 @@ class HexagonIOHandler { class HexagonRPCServer { public: - explicit HexagonRPCServer(uint8_t* receive_buffer, size_t receive_buffer_size_bytes) : io_{receive_buffer, receive_buffer_size_bytes}, rpc_server_{&io_} {}; + explicit HexagonRPCServer(uint8_t* receive_buffer, size_t receive_buffer_size_bytes) + : io_{receive_buffer, receive_buffer_size_bytes}, rpc_server_{&io_} {}; /*! * \brief Wrtie to IOHandler. @@ -179,11 +181,12 @@ class HexagonRPCServer { } // namespace tvm namespace { - tvm::runtime::hexagon::HexagonRPCServer* get_hexagon_rpc_server() { - static tvm::runtime::hexagon::HexagonRPCServer g_hexagon_rpc_server(new uint8_t[TVM_HEXAGON_RPC_BUFF_SIZE_BYTES], TVM_HEXAGON_RPC_BUFF_SIZE_BYTES); - return &g_hexagon_rpc_server; - } +tvm::runtime::hexagon::HexagonRPCServer* get_hexagon_rpc_server() { + static tvm::runtime::hexagon::HexagonRPCServer g_hexagon_rpc_server( + new uint8_t[TVM_HEXAGON_RPC_BUFF_SIZE_BYTES], TVM_HEXAGON_RPC_BUFF_SIZE_BYTES); + return &g_hexagon_rpc_server; } +} // namespace const tvm::runtime::PackedFunc get_runtime_func(const std::string& name) { if (const tvm::runtime::PackedFunc* pf = tvm::runtime::Registry::Get(name)) { @@ -227,7 +230,7 @@ int __QAIC_HEADER(hexagon_rpc_close)(remote_handle64 handle) { AEEResult __QAIC_HEADER(hexagon_rpc_send)(remote_handle64 _handle, const unsigned char* data, int dataLen) { int64_t written_size = get_hexagon_rpc_server()->Write(reinterpret_cast(data), - static_cast(dataLen)); + static_cast(dataLen)); if (written_size != dataLen) { HEXAGON_PRINT(ERROR, "RPC Server Write failed, written_size (%d) != dataLen (%d)", written_size, dataLen); From 532bf1bf4ff2b6955eba929ccabb64bd1370c1fa Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Tue, 7 Dec 2021 17:47:56 -0800 Subject: [PATCH 15/15] Fix build for other targets --- cmake/modules/Hexagon.cmake | 3 +++ src/runtime/hexagon/rpc/hexagon/rpc_server.cc | 1 + 2 files changed, 4 insertions(+) diff --git a/cmake/modules/Hexagon.cmake b/cmake/modules/Hexagon.cmake index 7d33b6b58d53..c1b2c3535b4f 100644 --- a/cmake/modules/Hexagon.cmake +++ b/cmake/modules/Hexagon.cmake @@ -75,6 +75,9 @@ if (NOT BUILD_FOR_HEXAGON AND NOT BUILD_FOR_ANDROID) USE_HEXAGON_PROXY_RPC STREQUAL "OFF" AND NOT USE_HEXAGON_RPC) if(USE_HEXAGON_DEVICE STREQUAL "OFF") list(APPEND COMPILER_SRCS src/target/opt/build_hexagon_off.cc) + if (NOT USE_HEXAGON_RPC) + return() + endif() elseif(NOT USE_HEXAGON_DEVICE STREQUAL "${PICK_SIM}" AND NOT USE_HEXAGON_DEVICE STREQUAL "${PICK_HW}") set(ERROR_MSG diff --git a/src/runtime/hexagon/rpc/hexagon/rpc_server.cc b/src/runtime/hexagon/rpc/hexagon/rpc_server.cc index bf4b0e3f5d85..f09223e1ad86 100644 --- a/src/runtime/hexagon/rpc/hexagon/rpc_server.cc +++ b/src/runtime/hexagon/rpc/hexagon/rpc_server.cc @@ -40,6 +40,7 @@ extern "C" { #include "../../hexagon/hexagon_common.h" #include "hexagon_rpc.h" +// TODO(mehrdadh): make this configurable. #define TVM_HEXAGON_RPC_BUFF_SIZE_BYTES 2 * 1024 * 1024 #define TVM_LOG_CUSTOMIZE 1