From 9b681833e4e0cfe815a0cb21d97b23f8defee4d0 Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Tue, 25 Jan 2022 16:42:24 -0800 Subject: [PATCH 1/6] Fix hexagon api build and Update Readme --- apps/hexagon_api/CMakeLists.txt | 3 ++ cmake/modules/Hexagon.cmake | 11 +------ python/tvm/contrib/hexagon/build.py | 2 +- .../contrib/test_hexagon/rpc/test_launcher.md | 32 +++++++++++++++---- 4 files changed, 30 insertions(+), 18 deletions(-) diff --git a/apps/hexagon_api/CMakeLists.txt b/apps/hexagon_api/CMakeLists.txt index 3c5eb616f1da..43edf589d961 100644 --- a/apps/hexagon_api/CMakeLists.txt +++ b/apps/hexagon_api/CMakeLists.txt @@ -88,3 +88,6 @@ ExternalProject_Add_Step(hexagon_tvm_runtime_rpc copy_binaries DEPENDEES install ) +# Copy android_bash template file +configure_file("${TVM_SOURCE_DIR}/src/runtime/hexagon/rpc/android_bash.sh.template" + ${HEXAGON_API_BINARY_DIR} COPYONLY) diff --git a/cmake/modules/Hexagon.cmake b/cmake/modules/Hexagon.cmake index d4dfaf22d698..1c7344c6157f 100644 --- a/cmake/modules/Hexagon.cmake +++ b/cmake/modules/Hexagon.cmake @@ -34,7 +34,7 @@ function(find_hexagon_toolchain) set(TRY_PATH "${USE_HEXAGON_SDK}") endif() message(STATUS "Looking for Hexagon toolchain in ${TRY_PATH}") - tvm_file_glob(GLOB_RECURSE HEXAGON_CLANG "${TRY_PATH}/*/hexagon-clang++") + file(GLOB_RECURSE HEXAGON_CLANG "${TRY_PATH}/*/hexagon-clang++") if(HEXAGON_CLANG) # The path is ${HEXAGON_TOOLCHAIN}/bin/hexagon-clang++. get_filename_component(HEXAGON_TMP0 "${HEXAGON_CLANG}" DIRECTORY) @@ -105,9 +105,6 @@ endif() # find_hexagon_sdk_root has been called at this point. if(USE_HEXAGON_RPC) - set(HEXAGON_RPC_OUTPUT "${CMAKE_CURRENT_BINARY_DIR}/hexagon_rpc") - file(MAKE_DIRECTORY ${HEXAGON_RPC_OUTPUT}) - set(TVMRT_SOURCE_DIR "${CMAKE_SOURCE_DIR}/src/runtime") set(QAIC_EXE "${HEXAGON_QAIC_EXE}") foreach(INCDIR IN LISTS HEXAGON_SDK_INCLUDES HEXAGON_REMOTE_ROOT) @@ -131,10 +128,6 @@ if(USE_HEXAGON_RPC) tvm_file_glob(GLOB RUNTIME_HEXAGON_SRCS "${TVMRT_SOURCE_DIR}/hexagon/rpc/android/*.cc") list(APPEND RUNTIME_HEXAGON_SRCS "${TVMRT_SOURCE_DIR}/hexagon/rpc/hexagon_rpc_stub.c") - # copy android_bash template file - configure_file("${TVMRT_SOURCE_DIR}/hexagon/rpc/android_bash.sh.template" - ${HEXAGON_RPC_OUTPUT} COPYONLY) - elseif(BUILD_FOR_HEXAGON) # Hexagon part find_hexagon_toolchain() @@ -154,8 +147,6 @@ if(USE_HEXAGON_RPC) SYSTEM PRIVATE "${TVMRT_SOURCE_DIR}/hexagon/rpc" ) endif() - - set_directory_properties(PROPERTIES ADDITIONAL_MAKE_CLEAN_FILES "${HEXAGON_RPC_OUTPUT}") endif() if(USE_HEXAGON_DEVICE STREQUAL "${PICK_SIM}") diff --git a/python/tvm/contrib/hexagon/build.py b/python/tvm/contrib/hexagon/build.py index e640aad89231..def9ea17ace0 100644 --- a/python/tvm/contrib/hexagon/build.py +++ b/python/tvm/contrib/hexagon/build.py @@ -49,7 +49,7 @@ def get_hexagon_rpc_dir() -> pathlib.Path: 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") + rpc_dir = os.path.join(os.path.dirname(path), "hexagon_api_output") if os.path.isdir(rpc_dir): HEXAGON_RPC_DIR = rpc_dir break diff --git a/tests/python/contrib/test_hexagon/rpc/test_launcher.md b/tests/python/contrib/test_hexagon/rpc/test_launcher.md index 463b88e3f374..13b8a8b11cea 100644 --- a/tests/python/contrib/test_hexagon/rpc/test_launcher.md +++ b/tests/python/contrib/test_hexagon/rpc/test_launcher.md @@ -26,20 +26,38 @@ Here are the steps that are taken to prepare a runtime on a Hexagon device to te - 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. +Note: before moving forward make sure to export Clang libraries to `LD_LIBRARY_PATH` and Hexagon toolchain to `HEXAGON_TOOLCHAIN`. + +To build these pieces, first build Hexagon API application under `apps/hexagon_api`. ```bash -cmake -DUSE_HEXAGON_RPC=ON \ - -DUSE_ANDROID_TOOLCHAIN=/path/to/android-ndk/build/cmake/android.toolchain.cmake \ +cd apps/hexagon_api +mkdir build +cd build +cmake -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_HEXAGON_TOOLCHAIN=/path/to/Hexagon/toolchain/"Tools"/sub-directory \ + -DUSE_OUTPUT_BINARY_DIR=/path/to/"tvm/build/hexagon_api_output" .. +``` + +This command generates `tvm_rpc_android` and `libtvm_runtime.so` to run on Android. Also, it generates `libtvm_runtime.a` and `libhexagon_rpc_skel.so` to run on Hexagon device. Now we have TVM artifacts which are used to run on the remote device. + +Next, we need to build TVM on host with RPC and Hexagon dependencies. To do that follow these commands. + +```bash +cd tvm +mkdir build +cd build +cmake -DUSE_LLVM=/path/to/llvm/bin/llvm-config \ -DUSE_CPP_RPC=ON \ - -DCMAKE_CXX_COMPILER=/path/to/clang++ \ - -DCMAKE_CXX_FLAGS='-stdlib=libc++' .. + -DCMAKE_CXX_COMPILER=/path/to/clang++ \ + -DCMAKE_CXX_FLAGS='-stdlib=libc++' \ + -DUSE_HEXAGON_SDK=/path/to/Hexagon/SDK \ + -DUSE_HEXAGON_ARCH=v65|v66|v68 \ + -DUSE_HEXAGON_DEVICE=target .. ``` ## Testing Using HexagonLauncher From d30e6ccb027420026f2535ef9484df3d8cf9f236 Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Tue, 25 Jan 2022 16:44:44 -0800 Subject: [PATCH 2/6] Cleanup hexagon_proxy_rpc --- apps/hexagon_proxy_rpc/Readme.md | 82 ----- apps/hexagon_proxy_rpc/cmake/HexagonRPC.cmake | 57 --- .../cmake/android/CMakeLists.txt | 104 ------ .../cmake/hexagon/CMakeLists.txt | 81 ----- apps/hexagon_proxy_rpc/common.h | 59 ---- apps/hexagon_proxy_rpc/hexagon_core.cc | 204 ----------- apps/hexagon_proxy_rpc/hexagon_proxy_rpc.idl | 35 -- apps/hexagon_proxy_rpc/rpc_env.cc | 326 ------------------ cmake/modules/Hexagon.cmake | 4 - 9 files changed, 952 deletions(-) delete mode 100644 apps/hexagon_proxy_rpc/Readme.md delete mode 100644 apps/hexagon_proxy_rpc/cmake/HexagonRPC.cmake delete mode 100644 apps/hexagon_proxy_rpc/cmake/android/CMakeLists.txt delete mode 100644 apps/hexagon_proxy_rpc/cmake/hexagon/CMakeLists.txt delete mode 100644 apps/hexagon_proxy_rpc/common.h delete mode 100644 apps/hexagon_proxy_rpc/hexagon_core.cc delete mode 100644 apps/hexagon_proxy_rpc/hexagon_proxy_rpc.idl delete mode 100644 apps/hexagon_proxy_rpc/rpc_env.cc diff --git a/apps/hexagon_proxy_rpc/Readme.md b/apps/hexagon_proxy_rpc/Readme.md deleted file mode 100644 index d7b577b2b378..000000000000 --- a/apps/hexagon_proxy_rpc/Readme.md +++ /dev/null @@ -1,82 +0,0 @@ - - - - - - - - - - - - - - - - -# Hexagon Proxy RPC server - -The proxy RPC server for Hexagon is a wrapper which takes standard TVM RPC calls from a python host -to a remote Android device and forwards them across FastRPC to Hexagon. This RPC flow will be replaced -by running a minimal RPC server directly on Hexagon. For now we provide a prototype forwarding RPC server -for host driven execution on Hexagon. - -## Compilation - -Project inventory: -* Android - * libtvm_runtime.so (containing HexagonHostDeviceAPI src/runtime/Hexagon/proxy_rpc/device_api.cc) - * tvm_rpc (C++ RPC server) - * librpc_env (Hexagon specific RPC proxy environment) - -* Hexagon - * libhexagon_proxy_rpc_skel.so (Hexagon device code containing FastRPC endpoints for the Hexagon Proxy RPC server) - -All Android and Hexagon device artifacts will be placed in `apps_hexagon_proxy_rpc` from which they can be pushed -to an attached `adb` device. - -### Prerequisites - -1. Android NDK version r19c or later. -2. Hexagon SDK version 4.0.0 or later. - -Android NDK can be downloaded from https://developer.android.com/ndk. -Hexagon SDK is available at //developer.qualcomm.com/software/Hexagon-dsp-sdk. - -### Compilation with TVM - -Building the Hexagon Proxy RPC as a component of the main TVM build -used for Hexagon codegen can be achieved by setting `USE_HEXAGON_PROXY_RPC=ON`. -A minimal example invocation for compiling TVM along with the Hexagon Proxy RPC server -is included below: - -``` -cmake -DCMAKE_C_COMPILER=/path/to/clang \ - -DCMAKE_CXX_COMPILER=/path/to/clang++ \ - -DCMAKE_CXX_FLAGS='-stdlib=libc++' \ - -DCMAKE_CXX_STANDARD=14 \ - -DUSE_RPC=ON \ - -DUSE_LLVM=/path/to/llvm/bin/llvm-config \ - -DUSE_HEXAGON_PROXY_RPC=ON \ - -DANDROID_ABI=arm64-v8a \ - -DANDROID_PLATFORM=android-28 \ - -DUSE_ANDROID_TOOLCHAIN=/path/to/android-ndk/build/cmake/android.toolchain.cmake \ - -DUSE_HEXAGON_ARCH=v65|v66|v68 \ - -DUSE_HEXAGON_SDK=/path/to/Hexagon/SDK \ - -DUSE_HEXAGON_TOOLCHAIN=/path/to/Hexagon/toolchain/ .. -``` - -where `v65|v66|v68` means "one of" these architecture versions. -The Hexagon proxy RPC application (tvm_rpc) is an android binary and thus requires the use -of an android toolchain for compilation. Similarly, the Hexagon tvm runtime -requires the use of the Hexagon toolchain and depends on the Hexagon SDK. The -resulting Hexagon launcher binaries can be found in the `apps_Hexagon_launcher` -subdirectory of the cmake build directory. The above command -will build support for Hexagon codegen in the TVM library that requires -`USE_LLVM` to be set to an llvm-config that has the Hexagon target built in. - - -# Disclaimer - -The Hexagon proxy RPC is intended for use with prototyping and does not utilize any -performance acceleration, as such the measured performance may be very poor. diff --git a/apps/hexagon_proxy_rpc/cmake/HexagonRPC.cmake b/apps/hexagon_proxy_rpc/cmake/HexagonRPC.cmake deleted file mode 100644 index 3ae6c8a7e664..000000000000 --- a/apps/hexagon_proxy_rpc/cmake/HexagonRPC.cmake +++ /dev/null @@ -1,57 +0,0 @@ -# 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. - -if(NOT DEFINED USE_HEXAGON_SDK) - message(SEND_ERROR "Please set USE_HEXAGON_SDK to the location of Hexagon SDK") -endif() -if (NOT DEFINED USE_HEXAGON_ARCH) - message(SEND_ERROR "Please set USE_HEXAGON_ARCH to the Hexagon architecture version") -endif() - -set(TVM_SOURCE_DIR "${CMAKE_CURRENT_SOURCE_DIR}/../../../../") - -include(ExternalProject) -include("${TVM_SOURCE_DIR}/cmake/utils/Utils.cmake") -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(QAIC_EXE "${HEXAGON_QAIC_EXE}") -foreach(INCDIR IN LISTS HEXAGON_SDK_INCLUDES HEXAGON_REMOTE_ROOT) - list(APPEND QAIC_FLAGS "-I${INCDIR}") -endforeach() - -set(HEXAGON_PROXY_RPC_SRC "${CMAKE_CURRENT_SOURCE_DIR}/../../") -set(CMAKE_SKIP_RPATH TRUE) - -# Qaic for the domain header. -# -# Don't add paths to these filenames, or otherwise cmake may spontaneously -# add -o option to the qaic invocation (with an undesirable path). -set(HEXAGON_PROXY_RPC_IDL "hexagon_proxy_rpc.idl") -set(HEXAGON_PROXY_RPC_H "hexagon_proxy_rpc.h") -set(HEXAGON_PROXY_RPC_SKEL_C "hexagon_proxy_rpc_skel.c") -set(HEXAGON_PROXY_RPC_STUB_C "hexagon_proxy_rpc_stub.c") - -include_directories( - "${HEXAGON_PROXY_RPC_SRC}" - "${TVM_SOURCE_DIR}/include" - "${TVM_SOURCE_DIR}/3rdparty/dlpack/include" - "${TVM_SOURCE_DIR}/3rdparty/dmlc-core/include" -) diff --git a/apps/hexagon_proxy_rpc/cmake/android/CMakeLists.txt b/apps/hexagon_proxy_rpc/cmake/android/CMakeLists.txt deleted file mode 100644 index 869456cce7e7..000000000000 --- a/apps/hexagon_proxy_rpc/cmake/android/CMakeLists.txt +++ /dev/null @@ -1,104 +0,0 @@ -# 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) -project(HexagonAndroidRPC C CXX) - -include("${CMAKE_CURRENT_SOURCE_DIR}/../HexagonRPC.cmake") - -add_custom_command( - OUTPUT ${HEXAGON_PROXY_RPC_STUB_C} ${HEXAGON_PROXY_RPC_H} - COMMAND ${QAIC_EXE} ${QAIC_FLAGS} "${HEXAGON_PROXY_RPC_SRC}/${HEXAGON_PROXY_RPC_IDL}" - MAIN_DEPENDENCY "${HEXAGON_PROXY_RPC_SRC}/${HEXAGON_PROXY_RPC_IDL}" -) - -include_directories(SYSTEM - "${HEXAGON_SDK_INCLUDES}" - "${HEXAGON_RPCMEM_ROOT}/inc" - "${CMAKE_CURRENT_BINARY_DIR}" # Output of qaic will go here -) - -link_directories(${HEXAGON_REMOTE_ROOT}) - -add_definitions(-DDMLC_USE_LOGGING_LIBRARY=) - -set(TVM_RPC_ENV_SOURCES - ${HEXAGON_PROXY_RPC_SRC}/rpc_env.cc -) - -add_library(rpc_env SHARED - ${TVM_RPC_ENV_SOURCES} - ${HEXAGON_PROXY_RPC_H} - ${HEXAGON_PROXY_RPC_STUB_C} -) - -ExternalProject_Add(android_tvm_runtime - SOURCE_DIR "${TVM_SOURCE_DIR}" - BUILD_COMMAND $(MAKE) runtime - CMAKE_ARGS - "-DCMAKE_TOOLCHAIN_FILE=${CMAKE_TOOLCHAIN_FILE}" - "-DANDROID_PLATFORM=${ANDROID_PLATFORM}" - "-DANDROID_ABI=${ANDROID_ABI}" - "-DCMAKE_CXX_STANDARD=14" - "-DUSE_LIBBACKTRACE=OFF" - "-DUSE_LLVM=OFF" - "-DUSE_RPC=ON" - "-DUSE_HEXAGON_SDK=${USE_HEXAGON_SDK}" - "-DUSE_HEXAGON_ARCH=${USE_HEXAGON_ARCH}" - INSTALL_COMMAND "" - BUILD_ALWAYS ON -) -ExternalProject_Get_Property(android_tvm_runtime BINARY_DIR) -ExternalProject_Add_Step(android_tvm_runtime copy_binaries - COMMAND ${CMAKE_COMMAND} -E copy_if_different - ${BINARY_DIR}/libtvm_runtime.so - ${CMAKE_CURRENT_BINARY_DIR} - DEPENDEES install -) - -add_dependencies(rpc_env android_tvm_runtime) -add_library(a_tvm_runtime SHARED IMPORTED) -set_target_properties(a_tvm_runtime PROPERTIES IMPORTED_LOCATION "${BINARY_DIR}/libtvm_runtime.so") - -target_link_libraries(rpc_env cdsprpc log a_tvm_runtime) - -# TVM CPP RPC build -set(TVM_RPC_DIR "${CMAKE_CURRENT_SOURCE_DIR}/../../../cpp_rpc") - - -set(TVM_RPC_SOURCES - ${TVM_RPC_DIR}/main.cc - ${TVM_RPC_DIR}/rpc_server.cc -) - -# Set output to same directory as the other TVM libs -set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}) -add_executable(tvm_rpc ${TVM_RPC_SOURCES}) - - -target_include_directories( - tvm_rpc - PUBLIC "${TVM_RPC_DIR}../../include" - PUBLIC "${TVM_RPC_DIR}../../3rdparty/dlpack" - PUBLIC "${TVM_RPC_DIR}../../3rdparty/dmlc-core" -) - -add_dependencies(rpc_env android_tvm_runtime) -target_link_libraries(rpc_env a_tvm_runtime) - -add_dependencies(tvm_rpc android_tvm_runtime rpc_env) -target_link_libraries(tvm_rpc a_tvm_runtime rpc_env) diff --git a/apps/hexagon_proxy_rpc/cmake/hexagon/CMakeLists.txt b/apps/hexagon_proxy_rpc/cmake/hexagon/CMakeLists.txt deleted file mode 100644 index 525212bab3b3..000000000000 --- a/apps/hexagon_proxy_rpc/cmake/hexagon/CMakeLists.txt +++ /dev/null @@ -1,81 +0,0 @@ -# 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) -project(HexagonRPCSkel C CXX) - -include("${CMAKE_CURRENT_SOURCE_DIR}/../HexagonRPC.cmake") - -add_custom_command( - OUTPUT ${HEXAGON_PROXY_RPC_SKEL_C} ${HEXAGON_PROXY_RPC_H} - COMMAND ${QAIC_EXE} ${QAIC_FLAGS} "${HEXAGON_PROXY_RPC_SRC}/${HEXAGON_PROXY_RPC_IDL}" - MAIN_DEPENDENCY "${HEXAGON_PROXY_RPC_SRC}/${HEXAGON_PROXY_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 - "${HEXAGON_PROXY_RPC_SRC}/hexagon_core.cc" -) - -add_library(hexagon_proxy_rpc_skel SHARED - "${HEXAGON_PROXY_RPC_H}" - "${HEXAGON_PROXY_RPC_SKEL_C}" - "${SKEL_SRCS}" -) - -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_proxy_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_proxy_rpc_skel -Wl,--whole-archive h_tvm_runtime -Wl,--no-whole-archive) diff --git a/apps/hexagon_proxy_rpc/common.h b/apps/hexagon_proxy_rpc/common.h deleted file mode 100644 index d93c90a6278c..000000000000 --- a/apps/hexagon_proxy_rpc/common.h +++ /dev/null @@ -1,59 +0,0 @@ -/* - * 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. - */ - -#ifndef TVM_RUNTIME_HEXAGON_PROXY_RPC_COMMON_H_ -#define TVM_RUNTIME_HEXAGON_PROXY_RPC_COMMON_H_ - -#include -#include -#include -#include -#include -#include - -#include -#include - -struct HandlePacket { - int ndim; - uint32_t handles[]; - int size() const { return size(ndim); } - static int size(int ndim) { return sizeof(HandlePacket) + ndim * sizeof(uint32_t); } -}; - -struct tensor_meta { - int ndim; - DLDataType dtype; - int64_t shape[]; - - int meta_size() const { return meta_size(ndim); } - int data_size() const { - int size = tvm::runtime::DataType(dtype).bytes(); - for (int d = 0; d != ndim; ++d) { - size *= shape[d]; - } - return size; - } - - static int meta_size(int ndim) { return sizeof(tensor_meta) + ndim * sizeof(int64_t); } - - std::string to_string() const; -}; - -#endif // TVM_RUNTIME_HEXAGON_PROXY_RPC_COMMON_H_ diff --git a/apps/hexagon_proxy_rpc/hexagon_core.cc b/apps/hexagon_proxy_rpc/hexagon_core.cc deleted file mode 100644 index e45bc24c30bf..000000000000 --- a/apps/hexagon_proxy_rpc/hexagon_core.cc +++ /dev/null @@ -1,204 +0,0 @@ -/* - * 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 "common.h" -#include "hexagon_proxy_rpc.h" - -template -T* DeserializeToPointerType(unsigned int module) { - return reinterpret_cast(module); -} - -template -unsigned int SerializeFromPointerType(T* pointer) { - return *reinterpret_cast(&pointer); -} - -tvm::runtime::Module load_module(const std::string& file_name) { - static const tvm::runtime::PackedFunc loader = - *tvm::runtime::Registry::Get("runtime.module.loadfile_hexagon"); - tvm::runtime::TVMRetValue rv = loader(file_name); - if (rv.type_code() == kTVMModuleHandle) { - return rv.operator tvm::runtime::Module(); - } - return tvm::runtime::Module(); -} - -int __QAIC_HEADER(hexagon_proxy_rpc_open)(const char* uri, remote_handle64* handle) { - FARF(ALWAYS, "[hexagon_proxy_rpc_open] FastRPC connection established"); - *handle = 0; - const tvm::runtime::PackedFunc api = *tvm::runtime::Registry::Get("device_api.hexagon.v2"); - tvm::runtime::Registry::Register("device_api.hexagon", true).set_body(api); - return AEE_SUCCESS; -} - -int __QAIC_HEADER(hexagon_proxy_rpc_close)(remote_handle64 handle) { - // Comment to stop clang-format from single-lining this function. - return AEE_SUCCESS; -} - -AEEResult __QAIC_HEADER(hexagon_proxy_rpc_load)(remote_handle64 handle, const char* module_path, - unsigned int* module) { - auto* mod_ptr = new tvm::runtime::Module(load_module(module_path)); - *module = SerializeFromPointerType(mod_ptr); - return AEE_SUCCESS; -} - -AEEResult __QAIC_HEADER(hexagon_proxy_rpc_unload)(remote_handle64 handle, unsigned int module) { - tvm::runtime::Module* mod_ptr = DeserializeToPointerType(module); - delete mod_ptr; - return AEE_SUCCESS; -} - -AEEResult __QAIC_HEADER(hexagon_proxy_rpc_get_function)(remote_handle64 handle, const char* name, - unsigned int module, unsigned int* func) { - tvm::runtime::Module* mod_ptr = DeserializeToPointerType(module); - std::string fname(name); - tvm::runtime::PackedFunc f = (*mod_ptr)->GetFunction(fname); - auto* f_ptr = new tvm::runtime::PackedFunc(f); - *func = SerializeFromPointerType(f_ptr); - return AEE_SUCCESS; -} - -AEEResult __QAIC_HEADER(hexagon_proxy_rpc_release_function)(remote_handle64 handle, - unsigned int func) { - tvm::runtime::PackedFunc* f_ptr = DeserializeToPointerType(func); - delete f_ptr; - return AEE_SUCCESS; -} - -AEEResult __QAIC_HEADER(hexagon_proxy_rpc_invoke)(remote_handle64 handle, unsigned int func, - const unsigned char* handles, int nhandles) { - tvm::runtime::PackedFunc* f_ptr = DeserializeToPointerType(func); - const auto* meta = reinterpret_cast(handles); - std::vector values; - std::vector type_codes; - for (size_t i = 0; i < meta->ndim; i++) { - tvm::runtime::NDArray* array = - DeserializeToPointerType(meta->handles[i]); - type_codes.push_back(kTVMDLTensorHandle); - values.emplace_back(); - const DLTensor* dltensor = array->operator->(); - values.back().v_handle = const_cast(static_cast(dltensor)); - } - - { - int res = qurt_hvx_reserve(QURT_HVX_RESERVE_ALL_AVAILABLE); - switch (res) { - case QURT_HVX_RESERVE_NOT_SUPPORTED: - case QURT_HVX_RESERVE_NOT_SUCCESSFUL: - FARF(ERROR, "error reserving HVX: %u", res); - return AEE_EFAILED; - default: - break; - } - // Lock HVX. - int lck = qurt_hvx_lock(QURT_HVX_MODE_128B); - if (lck != 0) { - FARF(ERROR, "error locking HVX: %u", lck); - return AEE_EFAILED; - } - } - tvm::runtime::TVMRetValue rv; - f_ptr->CallPacked(tvm::runtime::TVMArgs(values.data(), type_codes.data(), values.size()), &rv); - { - int unl = qurt_hvx_unlock(); - if (unl != 0) { - FARF(ERROR, "error unlocking HVX: %u", unl); - return AEE_EFAILED; - } - // Release HVX. - int rel = qurt_hvx_cancel_reserve(); - if (rel != 0) { - FARF(ERROR, "error canceling HVX reservation: %u", rel); - return AEE_EFAILED; - } - } - - return AEE_SUCCESS; -} - -AEEResult __QAIC_HEADER(hexagon_proxy_rpc_allocate)(remote_handle64 handle, - const unsigned char* input_meta, - int input_meta_size, const char* mem_scope, - unsigned int* tensor) { - const auto* meta = reinterpret_cast(input_meta); - auto device = tvm::Device{static_cast(kDLHexagon), 0}; - tvm::runtime::Optional scope; - if (*mem_scope) { - scope = mem_scope; - } - auto* array = new tvm::runtime::NDArray(std::move(tvm::runtime::NDArray::Empty( - tvm::ShapeTuple(meta->shape, meta->shape + meta->ndim), meta->dtype, device, scope))); - *tensor = SerializeFromPointerType(array); - return AEE_SUCCESS; -} - -AEEResult __QAIC_HEADER(hexagon_proxy_rpc_read)(remote_handle64 handle, unsigned char* dst_ptr, - int nbytes, unsigned int src) { - tvm::runtime::NDArray* src_ptr = DeserializeToPointerType(src); - const DLTensor* t = src_ptr->operator->(); - tvm::ShapeTuple shape(t->shape, t->shape + t->ndim); - auto* container = new tvm::runtime::NDArray::Container( - static_cast(dst_ptr), shape, src_ptr->operator->()->dtype, tvm::Device{kDLCPU, 0}); - container->SetDeleter([](tvm::Object* container) { - delete static_cast(container); - }); - tvm::runtime::NDArray dst(GetObjectPtr(container)); - dst.CopyFrom(*src_ptr); - return AEE_SUCCESS; -} - -AEEResult __QAIC_HEADER(hexagon_proxy_rpc_write)(remote_handle64 handle, unsigned int dst, - const unsigned char* src_ptr, int nbytes) { - tvm::runtime::NDArray* dst_ptr = DeserializeToPointerType(dst); - const DLTensor* t = dst_ptr->operator->(); - tvm::ShapeTuple shape(t->shape, t->shape + t->ndim); - auto* container = - new tvm::runtime::NDArray::Container(const_cast(src_ptr), shape, - dst_ptr->operator->()->dtype, tvm::Device{kDLCPU, 0}); - container->SetDeleter([](tvm::Object* container) { - delete static_cast(container); - }); - tvm::runtime::NDArray src(GetObjectPtr(container)); - dst_ptr->CopyFrom(src); - return AEE_SUCCESS; -} - -AEEResult __QAIC_HEADER(hexagon_proxy_rpc_release)(remote_handle64 handle, unsigned int array) { - tvm::runtime::NDArray* array_ptr = DeserializeToPointerType(array); - delete array_ptr; - return AEE_SUCCESS; -} diff --git a/apps/hexagon_proxy_rpc/hexagon_proxy_rpc.idl b/apps/hexagon_proxy_rpc/hexagon_proxy_rpc.idl deleted file mode 100644 index 0badf382d943..000000000000 --- a/apps/hexagon_proxy_rpc/hexagon_proxy_rpc.idl +++ /dev/null @@ -1,35 +0,0 @@ -/* - * 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_proxy_rpc : remote_handle64 { - AEEResult load(in string module_path, rout unsigned long mod); - AEEResult unload(in unsigned long mod); - AEEResult get_function(in string name, in unsigned long mod, rout unsigned long func); - AEEResult release_function(in unsigned long func); - AEEResult invoke(in unsigned long func, in buffer handles); - AEEResult allocate(in buffer template_tensor, in string mem_scope, rout unsigned long tensor); - AEEResult read(rout buffer dst_ptr, in unsigned long src); - AEEResult write(in unsigned long dst, in buffer src_ptr); - AEEResult release(in unsigned long array); -}; diff --git a/apps/hexagon_proxy_rpc/rpc_env.cc b/apps/hexagon_proxy_rpc/rpc_env.cc deleted file mode 100644 index 911ca580ba4f..000000000000 --- a/apps/hexagon_proxy_rpc/rpc_env.cc +++ /dev/null @@ -1,326 +0,0 @@ -/* - * 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 rpc_env.cc - * \brief Server environment of the RPC. - */ -#include "../cpp_rpc/rpc_env.h" - -#include -#include -#include -#include -#include -#include -#include -#include - -#include -#include -#include -#include - -#include "../../src/support/utils.h" -#include "common.h" -#include "hexagon_proxy_rpc.h" - -namespace tvm { -namespace runtime { - -/*! - * \brief CleanDir Removes the files from the directory - * \param dirname THe name of the directory - */ -void CleanDir(const std::string& dirname); - -namespace hexagon { -using FastRPCHandle = remote_handle64; -using Handle = uint32_t; - -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) { - std::cout << "error " << (enable ? "enabling" : "disabling") << " unsigned PD\n"; - } - return rc; -} - -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) { - std::cout << "error setting remote stack size: " << std::hex << rc << '\n'; - } - return rc; -} - -class FastRPCChannel { - public: - explicit FastRPCChannel(const std::string& uri) { - enable_unsigned_pd(true); - set_remote_stack_size(128 * 1024); - - int rc = hexagon_proxy_rpc_open(uri.c_str(), &handle_); - if (rc != AEE_SUCCESS) { - handle_ = std::numeric_limits::max(); - } - } - - ~FastRPCChannel() { - if (handle_ == std::numeric_limits::max()) { - return; - } - - hexagon_proxy_rpc_close(handle_); - handle_ = std::numeric_limits::max(); - } - - FastRPCHandle GetHandle() { return handle_; } - - private: - FastRPCHandle handle_ = std::numeric_limits::max(); -}; - -class HexagonModuleNode : public ModuleNode { - public: - HexagonModuleNode() = delete; - HexagonModuleNode(FastRPCHandle h, std::string file_name) : handle_(h), mod_{0} { - AEEResult rc = hexagon_proxy_rpc_load(handle_, file_name.c_str(), &mod_); - if (rc != AEE_SUCCESS) { - LOG(FATAL) << "Error loading module\n"; - } - } - ~HexagonModuleNode() { - AEEResult rc = hexagon_proxy_rpc_unload(handle_, mod_); - if (rc != AEE_SUCCESS) { - LOG(FATAL) << "Error unloading module\n"; - } - for (Handle func : packed_func_handles_) { - AEEResult rc = hexagon_proxy_rpc_release_function(handle_, func); - if (rc != AEE_SUCCESS) { - LOG(FATAL) << "Error releasing function\n"; - } - } - } - PackedFunc GetFunction(const std::string& name, const ObjectPtr& sptr_to_self) final { - hexagon::Handle func; - AEEResult rc = hexagon_proxy_rpc_get_function(handle_, name.c_str(), mod_, &func); - if (rc != AEE_SUCCESS) { - LOG(FATAL) << "Error calling get_function\n"; - } - packed_func_handles_.push_back(func); - return PackedFunc([handle = this->handle_, func, name](TVMArgs args, TVMRetValue* rv) { - std::vector handles; - for (size_t i = 0; i < args.size(); i++) { - ICHECK_EQ(args.type_codes[i], kTVMDLTensorHandle); - DLTensor* tensor = args[i]; - auto f = runtime::Registry::Get("runtime.hexagon.GetHandle"); - int32_t thandle = (*f)(tensor->data); - handles.push_back(thandle); - } - auto* packet = reinterpret_cast(rpcmem_alloc( - RPCMEM_HEAP_ID_SYSTEM, RPCMEM_DEFAULT_FLAGS, HandlePacket::size(args.size()))); - packet->ndim = args.size(); - std::copy(handles.begin(), handles.end(), packet->handles); - AEEResult rc = hexagon_proxy_rpc_invoke( - handle, func, reinterpret_cast(packet), packet->size()); - if (rc != AEE_SUCCESS) { - LOG(FATAL) << "Error invoking function: " << name; - } - rpcmem_free(packet); - }); - } - const char* type_key() const { return "HexagonModule"; } - - private: - FastRPCHandle handle_; - Handle mod_; - std::vector packed_func_handles_; -}; -} // namespace hexagon - -RPCEnv::RPCEnv(const std::string& wd) { - if (wd != "") { - base_ = wd + "/.cache"; - mkdir(wd.c_str(), 0777); - mkdir(base_.c_str(), 0777); - } else { - char cwd[PATH_MAX]; - auto cmdline = fopen("/proc/self/cmdline", "r"); - fread(cwd, 1, sizeof(cwd), cmdline); - fclose(cmdline); - std::string android_base_ = "/data/data/" + std::string(cwd) + "/cache"; - struct stat statbuf; - // Check if application data directory exist. If not exist, usually means we run tvm_rpc from - // adb shell terminal. - if (stat(android_base_.data(), &statbuf) == -1 || !S_ISDIR(statbuf.st_mode)) { - // Tmp directory is always writable for 'shell' user. - android_base_ = "/data/local/tmp"; - } - base_ = android_base_ + "/rpc"; - mkdir(base_.c_str(), 0777); - } - - static hexagon::FastRPCChannel hexagon_proxy_rpc(hexagon_proxy_rpc_URI CDSP_DOMAIN); - if (hexagon_proxy_rpc.GetHandle() == -1) { - LOG(FATAL) << "Error opening FastRPC channel\n"; - } - - TVM_REGISTER_GLOBAL("tvm.rpc.server.workpath").set_body([this](TVMArgs args, TVMRetValue* rv) { - *rv = this->GetPath(args[0]); - }); - - TVM_REGISTER_GLOBAL("tvm.rpc.server.load_module") - .set_body([this, handle = hexagon_proxy_rpc.GetHandle()](TVMArgs args, TVMRetValue* rv) { - std::string file_name = this->GetPath(args[0]); - auto n = make_object(handle, file_name); - *rv = Module(n); - LOG(INFO) << "Load module from " << file_name << " ..."; - }); - - TVM_REGISTER_GLOBAL("tvm.rpc.hexagon.allocate") - .set_body([handle = hexagon_proxy_rpc.GetHandle()](TVMArgs args, TVMRetValue* rv) { - DLTensor* ext_tensor = args[0]; - Optional mem_scope = args[1]; - - auto* input_meta = reinterpret_cast(rpcmem_alloc( - RPCMEM_HEAP_ID_SYSTEM, RPCMEM_DEFAULT_FLAGS, tensor_meta::meta_size(ext_tensor->ndim))); - input_meta->ndim = ext_tensor->ndim; - input_meta->dtype = ext_tensor->dtype; - std::copy(ext_tensor->shape, ext_tensor->shape + ext_tensor->ndim, input_meta->shape); - - hexagon::Handle hexagon_buffer; - const char* scope = mem_scope.defined() ? mem_scope.value().c_str() : ""; - AEEResult rc = - hexagon_proxy_rpc_allocate(handle, reinterpret_cast(input_meta), - input_meta->meta_size(), scope, &hexagon_buffer); - if (rc != AEE_SUCCESS) { - LOG(FATAL) << "Error allocating hexagon ndrray\n"; - } - rpcmem_free(input_meta); - *rv = static_cast(hexagon_buffer); - return rc == AEE_SUCCESS; - }); - - TVM_REGISTER_GLOBAL("tvm.rpc.hexagon.read_to_host") - .set_body([handle = hexagon_proxy_rpc.GetHandle()](TVMArgs args, TVMRetValue* rv) { - void* host_ptr = static_cast(args[0]); - size_t nbytes = args[1]; - hexagon::Handle hexagon_buffer = static_cast(args[2]); - AEEResult rc = hexagon_proxy_rpc_read(handle, static_cast(host_ptr), - static_cast(nbytes), hexagon_buffer); - if (rc != AEE_SUCCESS) { - LOG(FATAL) << "Error reading from hexagon buffer\n"; - } - }); - - TVM_REGISTER_GLOBAL("tvm.rpc.hexagon.write_from_host") - .set_body([handle = hexagon_proxy_rpc.GetHandle()](TVMArgs args, TVMRetValue* rv) { - hexagon::Handle hexagon_buffer = static_cast(args[0]); - void* host_ptr = static_cast(args[1]); - size_t nbytes = args[2]; - AEEResult rc = hexagon_proxy_rpc_write( - handle, hexagon_buffer, static_cast(host_ptr), static_cast(nbytes)); - if (rc != AEE_SUCCESS) { - LOG(FATAL) << "Error writing to hexagon buffer\n"; - } - }); - - TVM_REGISTER_GLOBAL("tvm.rpc.hexagon.release") - .set_body([handle = hexagon_proxy_rpc.GetHandle()](TVMArgs args, TVMRetValue* rv) { - hexagon::Handle hexagon_buffer = static_cast(args[0]); - AEEResult rc = hexagon_proxy_rpc_release(handle, hexagon_buffer); - if (rc != AEE_SUCCESS) { - LOG(FATAL) << "Error writing to hexagon buffer\n"; - } - }); -} - -/*! - * \brief GetPath To get the work path from packed function - * \param file_name The file name - * \return The full path of file. - */ -std::string RPCEnv::GetPath(const std::string& file_name) const { - // we assume file_name has "/" means file_name is the exact path - // and does not create /.rpc/ - return file_name.find('/') != std::string::npos ? file_name : base_ + "/" + file_name; -} -/*! - * \brief Remove The RPC Environment cleanup function - */ -void RPCEnv::CleanUp() const { - CleanDir(base_); - const int ret = rmdir(base_.c_str()); - if (ret != 0) { - LOG(WARNING) << "Remove directory " << base_ << " failed"; - } -} - -/*! - * \brief ListDir get the list of files in a directory - * \param dirname The root directory name - * \return vector Files in directory. - */ -std::vector ListDir(const std::string& dirname) { - std::vector vec; - DIR* dp = opendir(dirname.c_str()); - if (dp == nullptr) { - int errsv = errno; - LOG(FATAL) << "ListDir " << dirname << " error: " << strerror(errsv); - } - dirent* d; - while ((d = readdir(dp)) != nullptr) { - std::string filename = d->d_name; - if (filename != "." && filename != "..") { - std::string f = dirname; - if (f[f.length() - 1] != '/') { - f += '/'; - } - f += d->d_name; - vec.push_back(f); - } - } - closedir(dp); - return vec; -} - -/*! - * \brief CleanDir Removes the files from the directory - * \param dirname The name of the directory - */ -void CleanDir(const std::string& dirname) { - auto files = ListDir(dirname); - for (const auto& filename : files) { - std::string file_path = dirname + "/"; - file_path += filename; - const int ret = std::remove(filename.c_str()); - if (ret != 0) { - LOG(WARNING) << "Remove file " << filename << " failed"; - } - } -} -} // namespace runtime -} // namespace tvm diff --git a/cmake/modules/Hexagon.cmake b/cmake/modules/Hexagon.cmake index 1c7344c6157f..a990101bdecf 100644 --- a/cmake/modules/Hexagon.cmake +++ b/cmake/modules/Hexagon.cmake @@ -98,10 +98,6 @@ if(USE_HEXAGON_LAUNCHER STREQUAL "ON") message(SEND_ERROR "USE_HEXAGON_LAUNCHER is deprecated, please build apps separately") endif() -if(USE_HEXAGON_PROXY_RPC STREQUAL "ON") - message(SEND_ERROR "USE_HEXAGON_PROXY_RPC is deprecated, please build apps separately") -endif() - # find_hexagon_sdk_root has been called at this point. if(USE_HEXAGON_RPC) From ff2e9c7d853ced19b7cd4fafdc832c76203a0f5b Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Tue, 25 Jan 2022 16:56:24 -0800 Subject: [PATCH 3/6] Target Hack --- src/runtime/hexagon/android/hexagon_module.cc | 13 +++++++------ 1 file changed, 7 insertions(+), 6 deletions(-) diff --git a/src/runtime/hexagon/android/hexagon_module.cc b/src/runtime/hexagon/android/hexagon_module.cc index 3dc8fe9087c1..b2401b4e57bb 100644 --- a/src/runtime/hexagon/android/hexagon_module.cc +++ b/src/runtime/hexagon/android/hexagon_module.cc @@ -498,12 +498,13 @@ std::shared_ptr Device::Global() { #endif static std::shared_ptr dev( -#ifdef __ANDROID__ - CreateHexagonTarget() -#else - CreateHexagonSimulator() -#endif - ); // NOLINT + // TODO(mehrdadh): fix this hack + // #ifdef __ANDROID__ + // CreateHexagonTarget() + // #else + // CreateHexagonSimulator() + // #endif + nullptr); // NOLINT return dev; } From 172dd01840b2f763094f6497fd929eb23535583b Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Wed, 26 Jan 2022 08:30:06 -0800 Subject: [PATCH 4/6] Remove hack --- src/runtime/hexagon/android/hexagon_module.cc | 13 ++++++------- .../contrib/test_hexagon/rpc/test_launcher.md | 2 +- 2 files changed, 7 insertions(+), 8 deletions(-) diff --git a/src/runtime/hexagon/android/hexagon_module.cc b/src/runtime/hexagon/android/hexagon_module.cc index b2401b4e57bb..3dc8fe9087c1 100644 --- a/src/runtime/hexagon/android/hexagon_module.cc +++ b/src/runtime/hexagon/android/hexagon_module.cc @@ -498,13 +498,12 @@ std::shared_ptr Device::Global() { #endif static std::shared_ptr dev( - // TODO(mehrdadh): fix this hack - // #ifdef __ANDROID__ - // CreateHexagonTarget() - // #else - // CreateHexagonSimulator() - // #endif - nullptr); // NOLINT +#ifdef __ANDROID__ + CreateHexagonTarget() +#else + CreateHexagonSimulator() +#endif + ); // NOLINT return dev; } diff --git a/tests/python/contrib/test_hexagon/rpc/test_launcher.md b/tests/python/contrib/test_hexagon/rpc/test_launcher.md index 13b8a8b11cea..975902a15354 100644 --- a/tests/python/contrib/test_hexagon/rpc/test_launcher.md +++ b/tests/python/contrib/test_hexagon/rpc/test_launcher.md @@ -57,7 +57,7 @@ cmake -DUSE_LLVM=/path/to/llvm/bin/llvm-config \ -DCMAKE_CXX_FLAGS='-stdlib=libc++' \ -DUSE_HEXAGON_SDK=/path/to/Hexagon/SDK \ -DUSE_HEXAGON_ARCH=v65|v66|v68 \ - -DUSE_HEXAGON_DEVICE=target .. + -DUSE_HEXAGON_DEVICE=sim .. ``` ## Testing Using HexagonLauncher From 400b0ef6ce0ada9469d4217b04e8fd6d4b7efc41 Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Wed, 26 Jan 2022 16:27:45 -0800 Subject: [PATCH 5/6] address @cconvey comments --- apps/hexagon_api/CMakeLists.txt | 1 - .../contrib/test_hexagon/rpc/test_launcher.md | 29 ++++++++++--------- 2 files changed, 16 insertions(+), 14 deletions(-) diff --git a/apps/hexagon_api/CMakeLists.txt b/apps/hexagon_api/CMakeLists.txt index 43edf589d961..557dcfb85045 100644 --- a/apps/hexagon_api/CMakeLists.txt +++ b/apps/hexagon_api/CMakeLists.txt @@ -88,6 +88,5 @@ ExternalProject_Add_Step(hexagon_tvm_runtime_rpc copy_binaries DEPENDEES install ) -# Copy android_bash template file configure_file("${TVM_SOURCE_DIR}/src/runtime/hexagon/rpc/android_bash.sh.template" ${HEXAGON_API_BINARY_DIR} COPYONLY) diff --git a/tests/python/contrib/test_hexagon/rpc/test_launcher.md b/tests/python/contrib/test_hexagon/rpc/test_launcher.md index 975902a15354..bcf255e478f1 100644 --- a/tests/python/contrib/test_hexagon/rpc/test_launcher.md +++ b/tests/python/contrib/test_hexagon/rpc/test_launcher.md @@ -26,7 +26,13 @@ Here are the steps that are taken to prepare a runtime on a Hexagon device to te - Build TVM library with Hexagon support for host machine. - Build TVMRuntime library and C++ RPC server for host machine. -Note: before moving forward make sure to export Clang libraries to `LD_LIBRARY_PATH` and Hexagon toolchain to `HEXAGON_TOOLCHAIN`. +Note: First, ensure to export Clang libraries to `LD_LIBRARY_PATH` and Hexagon toolchain to `HEXAGON_TOOLCHAIN`: + +```bash +export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:"path to `llvm-clang/lib` sub-directory" + +export HEXAGON_TOOLCHAIN="Path to Hexagon toolchain. It can be the Hexagon toolchain included in the SDK, for example `HEXAGON_SDK_PATH/tools/HEXAGON_Tools/x.y.z/Tools`. The `x.y.z` in the path is the toolchain version number, which is specific to the version of the SDK." +``` To build these pieces, first build Hexagon API application under `apps/hexagon_api`. @@ -34,13 +40,13 @@ To build these pieces, first build Hexagon API application under `apps/hexagon_a cd apps/hexagon_api mkdir build cd build -cmake -DUSE_ANDROID_TOOLCHAIN=/path/to/android-ndk/build/cmake/android.toolchain.cmake \ +cmake -DUSE_ANDROID_TOOLCHAIN="path to `android-ndk/build/cmake/android.toolchain.cmake` file" \ -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/"Tools"/sub-directory \ - -DUSE_OUTPUT_BINARY_DIR=/path/to/"tvm/build/hexagon_api_output" .. + -DUSE_HEXAGON_SDK="path to Hexagon SDK" \ + -DUSE_HEXAGON_TOOLCHAIN="path to Hexagon toolchain `Tools` sub-directory which explained above" \ + -DUSE_OUTPUT_BINARY_DIR="path to `build/hexagon_api_output` which is a sub-directory of `tvm`" .. ``` This command generates `tvm_rpc_android` and `libtvm_runtime.so` to run on Android. Also, it generates `libtvm_runtime.a` and `libhexagon_rpc_skel.so` to run on Hexagon device. Now we have TVM artifacts which are used to run on the remote device. @@ -51,22 +57,19 @@ Next, we need to build TVM on host with RPC and Hexagon dependencies. To do that cd tvm mkdir build cd build -cmake -DUSE_LLVM=/path/to/llvm/bin/llvm-config \ +cmake -DUSE_LLVM="path to `llvm/bin/llvm-config`" \ -DUSE_CPP_RPC=ON \ - -DCMAKE_CXX_COMPILER=/path/to/clang++ \ + -DCMAKE_CXX_COMPILER="path to `clang++` executable" \ -DCMAKE_CXX_FLAGS='-stdlib=libc++' \ - -DUSE_HEXAGON_SDK=/path/to/Hexagon/SDK \ - -DUSE_HEXAGON_ARCH=v65|v66|v68 \ + -DUSE_HEXAGON_SDK="path to Hexagon SDK" \ + -DUSE_HEXAGON_ARCH="choose from v65|v66|v68" \ -DUSE_HEXAGON_DEVICE=sim .. ``` ## 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`. +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` as explained above. ```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 From 44e9a0779890c99162889f3529e9c78e90963410 Mon Sep 17 00:00:00 2001 From: Mehrdad Hessar Date: Wed, 26 Jan 2022 16:33:13 -0800 Subject: [PATCH 6/6] remove the rest of proxy rpc --- .../test_hexagon/proxy_rpc/__init__.py | 18 ----- .../test_hexagon/proxy_rpc/test_matmul.py | 73 ------------------- 2 files changed, 91 deletions(-) delete mode 100644 tests/python/contrib/test_hexagon/proxy_rpc/__init__.py delete mode 100644 tests/python/contrib/test_hexagon/proxy_rpc/test_matmul.py diff --git a/tests/python/contrib/test_hexagon/proxy_rpc/__init__.py b/tests/python/contrib/test_hexagon/proxy_rpc/__init__.py deleted file mode 100644 index 5261dc9cf052..000000000000 --- a/tests/python/contrib/test_hexagon/proxy_rpc/__init__.py +++ /dev/null @@ -1,18 +0,0 @@ -# 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 Proxy RPC """ diff --git a/tests/python/contrib/test_hexagon/proxy_rpc/test_matmul.py b/tests/python/contrib/test_hexagon/proxy_rpc/test_matmul.py deleted file mode 100644 index 839fdc9bc29d..000000000000 --- a/tests/python/contrib/test_hexagon/proxy_rpc/test_matmul.py +++ /dev/null @@ -1,73 +0,0 @@ -# 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 os - -import tvm -import tvm.testing -from tvm import te -import tvm.contrib.hexagon.hexagon as hexagon -from tvm.contrib import utils -import numpy as np - -from ..conftest import requires_hexagon_toolchain, requires_rpc_tracker_and_android_key - - -@requires_rpc_tracker_and_android_key -@requires_hexagon_toolchain -class TestMatMul: - M = tvm.testing.parameter(32) - N = tvm.testing.parameter(32) - K = tvm.testing.parameter(32) - - def test_matmul(self, M, N, K, rpc_sess, remote_path): - 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) - mod = tvm.build(schedule, [X, Y, Z], target=target_hexagon, target_host=target_hexagon) - - temp = utils.tempdir() - dso_binary_path = temp.relpath(os.path.basename(remote_path)) - mod.save(dso_binary_path) - - rpc_sess.upload(dso_binary_path, target=remote_path) - - mod = rpc_sess.load_module(remote_path) - - 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) - - dev = rpc_sess.hexagon(0) - xt = tvm.nd.array(x, device=dev) - yt = tvm.nd.array(y, device=dev) - zt = tvm.nd.array(z, device=dev) - 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) - - tvm.testing.assert_allclose(zt.asnumpy(), ztcpu.asnumpy(), rtol=1e-4)