You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by jr...@apache.org on 2021/12/08 18:44:32 UTC
[tvm] branch main updated: [Hexagon] Add RPC Mechanism for Hexagon (#9631)
This is an automated email from the ASF dual-hosted git repository.
jroesch pushed a commit to branch main
in repository https://gitbox.apache.org/repos/asf/tvm.git
The following commit(s) were added to refs/heads/main by this push:
new cd2fa69 [Hexagon] Add RPC Mechanism for Hexagon (#9631)
cd2fa69 is described below
commit cd2fa69677516048e165e84a88c774dfb0ee65d1
Author: Mehrdad Hessar <mh...@octoml.ai>
AuthorDate: Wed Dec 8 10:43:01 2021 -0800
[Hexagon] Add RPC Mechanism for Hexagon (#9631)
* Add Hexagon RPC
* removed android remote and updated Readme
* Add check for workspace size
* Make libtvm_runtime consistent for Android
* Remove root access
* Fix some docstrings
* Make stack remote size as parameter
* add documentation
* Refactor test conftest
* clang format
* Decoupled USE_HEXAGON_RPC
* fix creation of test base directory on android
* Address global variable
* Fix format and Cleanup cmake
* Fix build for other targets
---
CMakeLists.txt | 1 +
apps/cpp_rpc/CMakeLists.txt | 12 +-
cmake/libs/hexagon_rpc_skel/CMakeLists.txt | 119 ++++++++
cmake/modules/Hexagon.cmake | 138 +++++++++-
.../tvm/contrib/hexagon/__init__.py | 11 +-
python/tvm/contrib/hexagon/build.py | 300 +++++++++++++++++++++
python/tvm/contrib/{ => hexagon}/hexagon.py | 12 +-
python/tvm/contrib/hexagon/session.py | 75 ++++++
src/runtime/hexagon/rpc/android/session.cc | 120 +++++++++
.../runtime/hexagon/rpc/android_bash.sh.template | 14 +-
src/runtime/hexagon/rpc/hexagon/rpc_server.cc | 271 +++++++++++++++++++
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 +
tests/python/contrib/test_hexagon/conftest.py | 105 ++++++++
.../contrib/test_hexagon/proxy_rpc/test_matmul.py | 92 +------
.../test_hexagon/{conftest.py => rpc/__init__.py} | 10 +-
.../contrib/test_hexagon/{ => rpc}/conftest.py | 17 +-
.../contrib/test_hexagon/rpc/test_launcher.md | 98 +++++++
.../contrib/test_hexagon/rpc/test_launcher.py | 213 +++++++++++++++
.../python/unittest/test_target_codegen_hexagon.py | 4 +-
22 files changed, 1509 insertions(+), 144 deletions(-)
diff --git a/CMakeLists.txt b/CMakeLists.txt
index abf9f4a..2a4dfba 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 ccac53f..9664489 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 0000000..a4756aa
--- /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=<tvm/runtime/logging.h>)
+
+# 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 4c12bed..43789dd 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,12 @@ 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()
+ 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
@@ -202,6 +203,103 @@ 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.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)
+
+ set_directory_properties(PROPERTIES ADDITIONAL_MAKE_CLEAN_FILES "${HEXAGON_RPC_OUTPUT}")
+endif()
+
if(USE_HEXAGON_DEVICE STREQUAL "${PICK_SIM}")
find_hexagon_toolchain()
message(STATUS "Hexagon toolchain: ${HEXAGON_TOOLCHAIN}")
@@ -227,6 +325,7 @@ elseif(USE_HEXAGON_DEVICE STREQUAL "${PICK_HW}")
${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.
@@ -241,10 +340,39 @@ if (USE_HEXAGON_DEVICE STREQUAL "${PICK_NONE}")
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)
+ 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})
+
+ 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/tests/python/contrib/test_hexagon/conftest.py b/python/tvm/contrib/hexagon/__init__.py
similarity index 77%
copy from tests/python/contrib/test_hexagon/conftest.py
copy to python/tvm/contrib/hexagon/__init__.py
index b3bd00a..e728216 100644
--- a/tests/python/contrib/test_hexagon/conftest.py
+++ b/python/tvm/contrib/hexagon/__init__.py
@@ -14,13 +14,4 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
-
-""" Hexagon testing fixtures used to deduce testing argument
- values from testing parameters """
-
-import tvm
-
-
-@tvm.testing.fixture
-def shape_nhwc(batch, in_channel, in_size):
- return (batch, in_size, in_size, in_channel)
+"""Hexagon APIs."""
diff --git a/python/tvm/contrib/hexagon/build.py b/python/tvm/contrib/hexagon/build.py
new file mode 100644
index 0000000..ef081f2
--- /dev/null
+++ b/python/tvm/contrib/hexagon/build.py
@@ -0,0 +1,300 @@
+# 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 ..._ffi import libinfo
+from .session import Session
+
+
+RPC_SERVER_FILES = ["tvm_rpc_android", "libtvm_runtime.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, workspace_size_gb: int = 1):
+ """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._mod = None
+ self._workspace = None
+ self._workspace_max_size_mb = workspace_size_gb * 1024
+
+ HEXAGON_REMOTE_DEVICE_KEY = "hexagon-dev"
+
+ 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
+ """
+ # 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()
+
+ 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 "<RPC_TRACKER_HOST>" in line:
+ line = line.replace("<RPC_TRACKER_HOST>", str(rpc_tracker_host))
+ if "<RPC_TRACKER_PORT>" in line:
+ line = line.replace("<RPC_TRACKER_PORT>", str(rpc_tracker_port))
+ if "<HEXAGON_REMOTE_DEVICE_KEY>" in line:
+ line = line.replace(
+ "<HEXAGON_REMOTE_DEVICE_KEY>", self.HEXAGON_REMOTE_DEVICE_KEY
+ )
+ if "<RPC_SERVER_PORT>" in line:
+ line = line.replace("<RPC_SERVER_PORT>", 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}"
+ subprocess.check_call(self._adb_device_sub_cmd + ["push", src_path, destination])
+
+ # 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 hexagon_session_setup(self, remote_kw: dict):
+ """Setup Hexagon RPC Session from host to Hexagon device.
+
+ Parameters
+ ----------
+ remote_kw : dict
+ RPC tracker configs.
+ """
+ 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):
+ """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)
+ dst_remote_path = f"{self._workspace}/{remote_filename}"
+ subprocess.check_call(self._adb_device_sub_cmd + ["push", src_path, dst_remote_path])
+
+ 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.
+ 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.
+
+ 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 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`"]
+ )
+
+ 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])
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 fe25616..35136a3 100644
--- a/python/tvm/contrib/hexagon.py
+++ b/python/tvm/contrib/hexagon/hexagon.py
@@ -22,7 +22,7 @@ import os
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 @@ hexagon_link_main = os.path.join( # pylint: disable=invalid-name
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 link_shared(so_name, objs, **kwargs):
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 0000000..c413c60
--- /dev/null
+++ b/python/tvm/contrib/hexagon/session.py
@@ -0,0 +1,75 @@
+# 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",
+ 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
+
+ def __enter__(self):
+ if self.device:
+ # Already initialized
+ 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._remote_stack_size_bytes,
+ ],
+ )
+ 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 0000000..760b84e
--- /dev/null
+++ b/src/runtime/hexagon/rpc/android/session.cc
@@ -0,0 +1,120 @@
+/*
+ * 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 <tvm/runtime/registry.h>
+
+extern "C" {
+#include <AEEStdDef.h>
+#include <AEEStdErr.h>
+#include <HAP_farf.h>
+#include <HAP_perf.h>
+}
+
+#include <tvm/runtime/logging.h>
+
+#include <string>
+
+#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, int remote_stack_size_bytes) {
+ if (_handle != AEE_EUNKNOWN) return;
+
+ enable_unsigned_pd(true);
+ 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();
+ }
+
+ 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<const unsigned char*>(data), static_cast<int>(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<unsigned char*>(data),
+ static_cast<int>(size), &written_size);
+ ICHECK(rc == AEE_SUCCESS) << "hexagon_rpc_receive failed: " << rc;
+ return static_cast<size_t>(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<int>(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];
+ int remote_stack_size_bytes = args[1];
+ HexagonTransportChannel* hexagon_channel =
+ new HexagonTransportChannel(hexagon_rpc_URI CDSP_DOMAIN, remote_stack_size_bytes);
+ std::unique_ptr<RPCChannel> 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/tests/python/contrib/test_hexagon/conftest.py b/src/runtime/hexagon/rpc/android_bash.sh.template
similarity index 77%
copy from tests/python/contrib/test_hexagon/conftest.py
copy to src/runtime/hexagon/rpc/android_bash.sh.template
index b3bd00a..7bf6d77 100644
--- a/tests/python/contrib/test_hexagon/conftest.py
+++ b/src/runtime/hexagon/rpc/android_bash.sh.template
@@ -1,3 +1,4 @@
+#!/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
@@ -15,12 +16,9 @@
# specific language governing permissions and limitations
# under the License.
-""" Hexagon testing fixtures used to deduce testing argument
- values from testing parameters """
+export LD_LIBRARY_PATH=.
+./tvm_rpc_android server --port=<RPC_SERVER_PORT> --tracker=<RPC_TRACKER_HOST>:<RPC_TRACKER_PORT> --key=<HEXAGON_REMOTE_DEVICE_KEY>&
+rpc_pid=$!
-import tvm
-
-
-@tvm.testing.fixture
-def shape_nhwc(batch, in_channel, in_size):
- return (batch, in_size, in_size, in_channel)
+rm -f rpc_pid.txt
+echo $rpc_pid >> 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 0000000..f09223e
--- /dev/null
+++ b/src/runtime/hexagon/rpc/hexagon/rpc_server.cc
@@ -0,0 +1,271 @@
+/*
+ * 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 <AEEStdDef.h>
+#include <AEEStdErr.h>
+#include <HAP_farf.h>
+#include <HAP_perf.h>
+#include <qurt_error.h>
+#include <qurt_hvx.h>
+}
+
+#include <dlfcn.h>
+#include <tvm/runtime/object.h>
+#include <tvm/runtime/packed_func.h>
+#include <tvm/runtime/registry.h>
+
+#include <algorithm>
+#include <memory>
+#include <string>
+
+#include "../../../library_module.h"
+#include "../../../minrpc/minrpc_server.h"
+#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
+
+namespace tvm {
+namespace runtime {
+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, 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) {}
+
+ 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);
+ size_t written_size = static_cast<size_t>(
+ write_buffer_.sputn(reinterpret_cast<const char*>(buf), write_len_bytes));
+ if (written_size != write_len_bytes) {
+ HEXAGON_PRINT(ALWAYS, "HexagonIOHandler written_size failed");
+ }
+ return (ssize_t)written_size;
+ }
+
+ void MessageDone() {}
+
+ ssize_t PosixRead(uint8_t* buf, size_t read_len_bytes) {
+ HEXAGON_PRINT(ALWAYS, "HexagonIOHandler PosixRead called, %d, %d", read_len_bytes,
+ read_buffer_index_);
+
+ uint32_t bytes_to_read = 0;
+ 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_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);
+ }
+ return (ssize_t)bytes_to_read;
+ }
+
+ /*!
+ * \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_: ",
+ data_size_bytes, read_buffer_index_);
+ if (data_size_bytes > read_buffer_size_bytes_) {
+ return AEE_EFAILED;
+ }
+ read_buffer_ = data;
+ read_buffer_index_ = data_size_bytes;
+ return AEE_SUCCESS;
+ }
+
+ /*!
+ * \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.
+ *
+ * \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",
+ read_size_bytes);
+ return write_buffer_.sgetn(reinterpret_cast<char*>(buf), read_size_bytes);
+ }
+
+ void Close() { HEXAGON_PRINT(ALWAYS, "HexagonIOHandler Close called"); }
+
+ void Exit(int code) { exit(code); }
+
+ private:
+ const uint8_t* read_buffer_;
+ uint32_t read_buffer_index_;
+ size_t read_buffer_size_bytes_;
+
+ std::stringbuf write_buffer_;
+};
+
+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_} {};
+
+ /*!
+ * \brief Wrtie to IOHandler.
+ * \param data The data pointer
+ * \param data_size_bytes The data size in bytes.
+ *
+ * \return The size of data written to IOHandler.
+ */
+ int64_t Write(const uint8_t* data, size_t data_size_bytes) {
+ if (io_.SetReadBuffer(data, data_size_bytes) != AEE_SUCCESS) {
+ return -1;
+ }
+ rpc_server_.ProcessOnePacket();
+ return (int64_t)data_size_bytes;
+ }
+
+ /*!
+ * \brief Read from IOHandler.
+ * \param buf The buffer pointer
+ * \param read_size_bytes Read request size 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);
+ }
+
+ private:
+ HexagonIOHandler io_;
+ MinRPCServer<HexagonIOHandler> rpc_server_;
+};
+
+} // namespace hexagon
+} // namespace runtime
+} // 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;
+}
+} // namespace
+
+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<remote_handle64>(reinterpret_cast<uintptr_t>(malloc(1)));
+ if (!*handle) {
+ HEXAGON_PRINT(ERROR, "%s: cannot allocate memory", __func__);
+ return AEE_ENOMEMORY;
+ }
+ reset_device_api();
+ get_hexagon_rpc_server();
+ return AEE_SUCCESS;
+}
+
+int __QAIC_HEADER(hexagon_rpc_close)(remote_handle64 handle) {
+ HEXAGON_PRINT(ALWAYS, "%s", __func__);
+ if (handle) {
+ free(reinterpret_cast<void*>(static_cast<uintptr_t>(handle)));
+ }
+ return AEE_SUCCESS;
+}
+
+/*!
+ * \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.
+ *
+ * \return The status.
+ */
+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<const uint8_t*>(data),
+ static_cast<size_t>(dataLen));
+ if (written_size != dataLen) {
+ HEXAGON_PRINT(ERROR, "RPC Server Write failed, written_size (%d) != dataLen (%d)", written_size,
+ 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.
+ *
+ * \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 =
+ get_hexagon_rpc_server()->Read(reinterpret_cast<uint8_t*>(buf), static_cast<size_t>(bufLen));
+ *buf_written_size = read_size;
+ if (read_size == bufLen) {
+ return AEE_SUCCESS;
+ } else {
+ HEXAGON_PRINT(ALWAYS, "RPC Server Read failed, read_size (%d) != dataLen (%d)", read_size,
+ bufLen);
+ 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<tvm::runtime::Library> 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 0000000..55b8d39
--- /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<octet> buffer;
+
+interface hexagon_rpc : remote_handle64 {
+ AEEResult send(in buffer data);
+ 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 0cc5c5b..fed2ad5 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<PrimExpr> 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 fcc44fb..f31f3f6 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 <llvm/Analysis/TargetTransformInfo.h>
@@ -72,11 +71,7 @@
#include <llvm/Support/FileSystem.h>
#include <llvm/Support/Host.h>
#include <llvm/Support/MemoryBuffer.h>
-#if TVM_LLVM_VERSION >= 140
-#include <llvm/MC/TargetRegistry.h>
-#else
#include <llvm/Support/TargetRegistry.h>
-#endif
#include <llvm/Support/TargetSelect.h>
#include <llvm/Support/raw_ostream.h>
#include <llvm/Target/TargetMachine.h>
diff --git a/tests/lint/check_file_type.py b/tests/lint/check_file_type.py
index 1b45ac7..9640038 100644
--- a/tests/lint/check_file_type.py
+++ b/tests/lint/check_file_type.py
@@ -150,6 +150,8 @@ ALLOW_SPECIFIC_FILE = {
"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/conftest.py b/tests/python/contrib/test_hexagon/conftest.py
index b3bd00a..bd2ae7c 100644
--- a/tests/python/contrib/test_hexagon/conftest.py
+++ b/tests/python/contrib/test_hexagon/conftest.py
@@ -18,9 +18,114 @@
""" 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 2b18911..839fdc9 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,18 @@
# under the License.
import os
-import sys
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
-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/conftest.py b/tests/python/contrib/test_hexagon/rpc/__init__.py
similarity index 77%
copy from tests/python/contrib/test_hexagon/conftest.py
copy to tests/python/contrib/test_hexagon/rpc/__init__.py
index b3bd00a..92e96bf 100644
--- a/tests/python/contrib/test_hexagon/conftest.py
+++ b/tests/python/contrib/test_hexagon/rpc/__init__.py
@@ -15,12 +15,4 @@
# specific language governing permissions and limitations
# under the License.
-""" Hexagon testing fixtures used to deduce testing argument
- values from testing parameters """
-
-import tvm
-
-
-@tvm.testing.fixture
-def shape_nhwc(batch, in_channel, in_size):
- return (batch, in_size, in_size, in_channel)
+""" Testing infrastructure for Hexagon RPC"""
diff --git a/tests/python/contrib/test_hexagon/conftest.py b/tests/python/contrib/test_hexagon/rpc/conftest.py
similarity index 72%
copy from tests/python/contrib/test_hexagon/conftest.py
copy to tests/python/contrib/test_hexagon/rpc/conftest.py
index b3bd00a..50c199a 100644
--- a/tests/python/contrib/test_hexagon/conftest.py
+++ b/tests/python/contrib/test_hexagon/rpc/conftest.py
@@ -15,12 +15,17 @@
# specific language governing permissions and limitations
# under the License.
-""" Hexagon testing fixtures used to deduce testing argument
- values from testing parameters """
+import pytest
-import tvm
+def pytest_addoption(parser):
+ parser.addoption(
+ "--serial-number",
+ required=True,
+ help=("Android device serial number list from 'adb' command."),
+ )
-@tvm.testing.fixture
-def shape_nhwc(batch, in_channel, in_size):
- return (batch, in_size, in_size, in_channel)
+
+@pytest.fixture
+def android_serial_number(request):
+ return request.config.getoption("--serial-number")
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 0000000..463b88e
--- /dev/null
+++ b/tests/python/contrib/test_hexagon/rpc/test_launcher.md
@@ -0,0 +1,98 @@
+<!--- 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. -->
+
+# 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. 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
+```
+
+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()
+
+# Create an RPC session from host to Hexagon.
+remote_kw = {
+ "host": "TVM_TRACKER_HOST",
+ "port": "TVM_TRACKER_PORT",
+ "priority": 0,
+ "timeout": 60,
+}
+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 0000000..d705541
--- /dev/null
+++ b/tests/python/contrib/test_hexagon/rpc/test_launcher.py
@@ -0,0 +1,213 @@
+# 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_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)
+ 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_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"
+ 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.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
+ @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")
+ 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], tvm.target.Target(target_hexagon, 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.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], tvm.target.Target(target_llvm, 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.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))
+ 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,
+ tvm.target.Target(target_hexagon, 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.hexagon_session_setup(remote_kw)
+ launcher.upload(dso_binary_path, 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)
+ 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,
+ tvm.target.Target(target_llvm, 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 220fa86..ef0eb4f 100644
--- a/tests/python/unittest/test_target_codegen_hexagon.py
+++ b/tests/python/unittest/test_target_codegen_hexagon.py
@@ -23,12 +23,12 @@ import sys
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