From 628c6a86b44581ed7e4e9371e9fb77aaa8fed2f8 Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Mon, 27 Jul 2020 16:07:29 +0100 Subject: [PATCH 01/37] [BYOC][ETHOSN] Introduce the Ethos-N BYOC integration This is the first of 3 PRs to introduce the Ethos-N integration into TVM via the BYOC framework. It adds support for partitioning and compiling for the Ethos-N77 target with CPU fallback for unsupported operators. Additionally, runtime support is added in the form of an Ethos-N runtime module. In this initial PR, only quantized concatenate and split are supported with follow-up PRs adding support for many further operators. Co-authored-by: Leo Blonk Co-authored-by: Tristan O'Connor Co-authored-by: Leandro Nunes Co-authored-by: Ramana Radhakrishnan Co-authored-by: Luke Hutton --- CMakeLists.txt | 7 + cmake/config.cmake | 10 + cmake/modules/contrib/EthosN.cmake | 54 +++ cmake/util/FindEthosN.cmake | 95 +++++ python/tvm/relay/op/contrib/__init__.py | 1 + python/tvm/relay/op/contrib/_ethosn.py | 22 ++ python/tvm/relay/op/contrib/ethosn.py | 64 ++++ .../backend/contrib/ethosn/capabilities.h | 57 +++ src/relay/backend/contrib/ethosn/codegen.cc | 214 +++++++++++ .../backend/contrib/ethosn/codegen_ethosn.h | 331 ++++++++++++++++++ .../backend/contrib/ethosn/ethosn_api.cc | 268 ++++++++++++++ src/relay/backend/contrib/ethosn/ethosn_api.h | 142 ++++++++ src/runtime/contrib/ethosn/ethosn_device.cc | 222 ++++++++++++ src/runtime/contrib/ethosn/ethosn_device.h | 44 +++ src/runtime/contrib/ethosn/ethosn_runtime.cc | 146 ++++++++ src/runtime/contrib/ethosn/ethosn_runtime.h | 110 ++++++ tests/python/contrib/test_ethosn/__init__.py | 18 + .../contrib/test_ethosn/_infrastructure.py | 22 ++ .../contrib/test_ethosn/infrastructure.py | 225 ++++++++++++ .../contrib/test_ethosn/test_concatenate.py | 90 +++++ .../python/contrib/test_ethosn/test_split.py | 70 ++++ .../contrib/test_ethosn/test_topologies.py | 122 +++++++ tests/scripts/task_config_build_cpu.sh | 2 + tests/scripts/task_python_ethosn_tests.sh | 30 ++ 24 files changed, 2366 insertions(+) create mode 100644 cmake/modules/contrib/EthosN.cmake create mode 100644 cmake/util/FindEthosN.cmake create mode 100644 python/tvm/relay/op/contrib/_ethosn.py create mode 100644 python/tvm/relay/op/contrib/ethosn.py create mode 100644 src/relay/backend/contrib/ethosn/capabilities.h create mode 100644 src/relay/backend/contrib/ethosn/codegen.cc create mode 100644 src/relay/backend/contrib/ethosn/codegen_ethosn.h create mode 100644 src/relay/backend/contrib/ethosn/ethosn_api.cc create mode 100644 src/relay/backend/contrib/ethosn/ethosn_api.h create mode 100644 src/runtime/contrib/ethosn/ethosn_device.cc create mode 100644 src/runtime/contrib/ethosn/ethosn_device.h create mode 100644 src/runtime/contrib/ethosn/ethosn_runtime.cc create mode 100644 src/runtime/contrib/ethosn/ethosn_runtime.h create mode 100644 tests/python/contrib/test_ethosn/__init__.py create mode 100644 tests/python/contrib/test_ethosn/_infrastructure.py create mode 100644 tests/python/contrib/test_ethosn/infrastructure.py create mode 100644 tests/python/contrib/test_ethosn/test_concatenate.py create mode 100644 tests/python/contrib/test_ethosn/test_split.py create mode 100644 tests/python/contrib/test_ethosn/test_topologies.py create mode 100755 tests/scripts/task_python_ethosn_tests.sh diff --git a/CMakeLists.txt b/CMakeLists.txt index 0565cfd939bb..e4d40a560afd 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -8,6 +8,7 @@ include(cmake/util/FindOpenCL.cmake) include(cmake/util/FindVulkan.cmake) include(cmake/util/FindLLVM.cmake) include(cmake/util/FindROCM.cmake) +include(cmake/util/FindEthosN.cmake) if(EXISTS ${CMAKE_CURRENT_BINARY_DIR}/config.cmake) include(${CMAKE_CURRENT_BINARY_DIR}/config.cmake) @@ -44,6 +45,7 @@ tvm_option(INSTALL_DEV "Install compiler infrastructure" OFF) tvm_option(HIDE_PRIVATE_SYMBOLS "Compile with -fvisibility=hidden." OFF) tvm_option(USE_TF_TVMDSOOP "Build with TensorFlow TVMDSOOp" OFF) tvm_option(USE_FALLBACK_STL_MAP "Use TVM's POD compatible Map" OFF) +tvm_option(USE_ETHOSN "Build with Arm Ethos-N" OFF) # 3rdparty libraries tvm_option(DLPACK_PATH "Path to DLPACK" "3rdparty/dlpack/include") @@ -307,6 +309,11 @@ include(cmake/modules/Metal.cmake) include(cmake/modules/ROCM.cmake) include(cmake/modules/LLVM.cmake) include(cmake/modules/Micro.cmake) +<<<<<<< HEAD +======= +include(cmake/modules/ANTLR.cmake) +include(cmake/modules/contrib/EthosN.cmake) +>>>>>>> [BYOC][ETHOSN] Introduce the Ethos-N BYOC integration include(cmake/modules/contrib/BLAS.cmake) include(cmake/modules/contrib/CODEGENC.cmake) include(cmake/modules/contrib/DNNL.cmake) diff --git a/cmake/config.cmake b/cmake/config.cmake index fb4e2bd3e277..3bd074d22851 100644 --- a/cmake/config.cmake +++ b/cmake/config.cmake @@ -210,6 +210,16 @@ set(USE_DNNL_CODEGEN OFF) set(USE_ARM_COMPUTE_LIB OFF) set(USE_ARM_COMPUTE_LIB_GRAPH_RUNTIME OFF) +# Whether to build with Arm Ethos-N support +# Possible values: +# - OFF: disable Arm Ethos-N support +# - path/to/arm-ethos-N-stack: use a specific version of the +# Ethos-N driver stack +set(USE_ETHOSN OFF) +# If USE_ETHOSN is enabled, use Ethos-N hardware (ON) or +# software test infrastructure (OFF) +set(USE_ETHOSN_HW ON) + # Build ANTLR parser for Relay text format # Possible values: # - ON: enable ANTLR by searching default locations (cmake find_program for antlr4 and /usr/local for jar) diff --git a/cmake/modules/contrib/EthosN.cmake b/cmake/modules/contrib/EthosN.cmake new file mode 100644 index 000000000000..685ea22bebcf --- /dev/null +++ b/cmake/modules/contrib/EthosN.cmake @@ -0,0 +1,54 @@ +# 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. + +# Arm Ethos-N rules + +if(NOT USE_ETHOSN STREQUAL "OFF") + find_ethosn(${USE_ETHOSN}) + + if(NOT ETHOSN_FOUND) + message(FATAL_ERROR "Cannot find Ethos-N, USE_ETHOSN=" ${USE_ETHOSN}) + endif() + + if (ETHOSN_FOUND) + include_directories(${ETHOSN_INCLUDE_DIRS}) + add_definitions(${ETHOSN_DEFINITIONS}) + + message(STATUS "Build with Ethos-N ${ETHOSN_PACKAGE_VERSION}") + + file(GLOB ETHOSN_RUNTIME_CONTRIB_SRC + CONFIGURE_DEPENDS src/runtime/contrib/ethosn/ethosn_runtime.cc + CONFIGURE_DEPENDS src/runtime/contrib/ethosn/ethosn_device.cc) + list(APPEND RUNTIME_SRCS ${ETHOSN_RUNTIME_CONTRIB_SRC}) + + file(GLOB COMPILER_ETHOSN_SRCS + CONFIGURE_DEPENDS src/relay/backend/contrib/ethosn/*) + list(APPEND COMPILER_SRCS ${COMPILER_ETHOSN_SRCS}) + + list(APPEND TVM_LINKER_LIBS ${ETHOSN_COMPILER_LIBRARY} + ${ETHOSN_RUNTIME_LIBRARY}) + list(APPEND TVM_RUNTIME_LINKER_LIBS ${ETHOSN_COMPILER_LIBRARY} + ${ETHOSN_RUNTIME_LIBRARY}) + + if(NOT MSVC) + set_source_files_properties(${COMPILER_ETHOSN_SRCS} + PROPERTIES COMPILE_DEFINITIONS "DMLC_ENABLE_RTTI=0") + set_source_files_properties(${COMPILER_ETHOSN_SRCS} + PROPERTIES COMPILE_FLAGS "-fno-rtti") + endif() + endif(ETHOSN_FOUND) +endif(NOT USE_ETHOSN STREQUAL "OFF") diff --git a/cmake/util/FindEthosN.cmake b/cmake/util/FindEthosN.cmake new file mode 100644 index 000000000000..7e8c7b90b756 --- /dev/null +++ b/cmake/util/FindEthosN.cmake @@ -0,0 +1,95 @@ +# 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. + +####################################################### +# Find Arm Ethos-N libraries +# +# Usage: +# find_ethosn(${USE_ETHOSN}) +# +# - When USE_ETHOSN=/path/to/ethos-sdk-path, use the path from USE_ETHOSN +# - Else, when environment variable ETHOSN_STACK is set, use that path +# - When USE_ETHOSN=ON, use auto search +# +# Provide variables: +# +# - ETHOSN_FOUND +# - ETHOSN_PACKAGE_VERSION +# - ETHOSN_DEFINITIONS +# - ETHOSN_INCLUDE_DIRS +# - ETHOSN_COMPILER_LIBRARY +# - ETHOSN_RUNTIME_LIBRARY + +macro(find_ethosn use_ethosn) + set(__use_ethosn ${use_ethosn}) + if(IS_DIRECTORY ${__use_ethosn}) + set(__ethosn_stack ${__use_ethosn}) + message(STATUS "Arm Ethos-N driver stack PATH=" ${__use_ethosn}) + elseif(IS_DIRECTORY $ENV{ETHOSN_STACK}) + set(__ethosn_stack $ENV{ETHOSN_STACK}) + message(STATUS "Arm Ethos-N driver stack from env=" ${__use_ethosn}) + else() + set(__ethosn_stack "") + endif() + + if(__ethosn_stack) + set(ETHOSN_INCLUDE_DIRS "") + # Compile-time support + find_path(_SL_DIR NAMES Support.hpp + PATHS ${__ethosn_stack}/include/ethosn_support_library) + string(REGEX REPLACE "/ethosn_support_library" "" _SL_DIR2 ${_SL_DIR}) + list(APPEND ETHOSN_INCLUDE_DIRS "${_SL_DIR2}") + + find_library(ETHOSN_COMPILER_LIBRARY NAMES EthosNSupport + PATHS ${__ethosn_stack}/lib) + find_library(ETHOSN_COMPILER_LIBRARY NAMES EthosNSupport) + + set(ETHOSN_PACKAGE_VERSION "0.1.1") + + if(USE_ETHOSN_HW STREQUAL "ON") + # Runtime hardware support + find_path(_DL_DIR NAMES Network.hpp + PATHS ${__ethosn_stack}/include/ethosn_driver_library) + string(REGEX REPLACE "/ethosn_driver_library" "" _DL_DIR2 ${_DL_DIR}) + list(APPEND ETHOSN_INCLUDE_DIRS "${_DL_DIR2}") + + find_library(ETHOSN_RUNTIME_LIBRARY NAMES EthosNDriver + PATHS ${__ethosn_stack}/lib) + find_library(ETHOSN_RUNTIME_LIBRARY NAMES EthosNDriver) + set(ETHOSN_DEFINITIONS -DETHOSN_HW) + endif () + + if(ETHOSN_COMPILER_LIBRARY) + set(ETHOSN_FOUND TRUE) + endif() + endif(__ethosn_stack) + + if(NOT ETHOSN_FOUND) + if(__use_ethosn STREQUAL "ON") + message(WARNING "No cmake find_package available for Arm Ethos-N") + endif() + endif() + + # additional libraries + if(ETHOSN_FOUND) + message(STATUS "Found ETHOSN_DEFINITIONS=${ETHOSN_DEFINITIONS}") + message(STATUS "Found ETHOSN_INCLUDE_DIRS=${ETHOSN_INCLUDE_DIRS}") + message(STATUS "Found ETHOSN_COMPILER_LIBRARY=${ETHOSN_COMPILER_LIBRARY}") + message(STATUS "Found ETHOSN_RUNTIME_LIBRARY=${ETHOSN_RUNTIME_LIBRARY}") + endif(ETHOSN_FOUND) + +endmacro(find_ethosn) diff --git a/python/tvm/relay/op/contrib/__init__.py b/python/tvm/relay/op/contrib/__init__.py index 26ca78c1190b..dbcd8055d30b 100644 --- a/python/tvm/relay/op/contrib/__init__.py +++ b/python/tvm/relay/op/contrib/__init__.py @@ -21,3 +21,4 @@ from .arm_compute_lib import * from .dnnl import * from .coreml import * +from .ethosn import * diff --git a/python/tvm/relay/op/contrib/_ethosn.py b/python/tvm/relay/op/contrib/_ethosn.py new file mode 100644 index 000000000000..ea2915675ec6 --- /dev/null +++ b/python/tvm/relay/op/contrib/_ethosn.py @@ -0,0 +1,22 @@ +# 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. + +"""Expose 'is supported' functions to Python.""" + +import tvm._ffi + +tvm._ffi._init_api("relay.ethos-n.support", __name__) diff --git a/python/tvm/relay/op/contrib/ethosn.py b/python/tvm/relay/op/contrib/ethosn.py new file mode 100644 index 000000000000..3f0121cd2b5e --- /dev/null +++ b/python/tvm/relay/op/contrib/ethosn.py @@ -0,0 +1,64 @@ +# 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. +# pylint: disable=invalid-name, unused-argument +"""Arm(R) Ethos(TM) -N NPU supported operators.""" +import tvm.ir +from ... import qnn as _qnn +from . import _ethosn as support + + +@tvm.ir.register_op_attr("qnn.concatenate", "target.ethos-n") +def qnn_concatenate(attrs, args): + """Check if a concatenate is supported by Ethos-N.""" + conc = _qnn.op.concatenate(*args, **attrs) + if not support.concatenate(conc): + return False + + # Support library has some unenforced restrictions on qnn params + min_range = 1e9 + max_range = -1e9 + qnn_params = [] + for i in range(len(args[1].fields)): + scale = args[1].fields[i].data.asnumpy() + zero_point = args[2].fields[i].data.asnumpy() + min_range = min(-1 * zero_point * scale, min_range) + max_range = max((255 - zero_point) * scale, max_range) + qnn_params.append((scale, zero_point)) + + scale = (max_range - min_range) / 255 + zero_point = int(-min_range/scale) + if (scale, zero_point) in qnn_params: + return True + + return False + + +@tvm.ir.register_op_attr("split", "target.ethos-n") +def split(attrs, args): + """Check if a split is supported by Ethos-N.""" + if isinstance(attrs["indices_or_sections"], tvm.tir.IntImm): + sp = tvm.relay.split(*args, + indices_or_sections=attrs["indices_or_sections"].value, + axis=attrs["axis"]) + else: + sp = tvm.relay.split(*args, + indices_or_sections=attrs["indices_or_sections"], + axis=attrs["axis"]) + if not support.split(sp.astuple()): + return False + + return True diff --git a/src/relay/backend/contrib/ethosn/capabilities.h b/src/relay/backend/contrib/ethosn/capabilities.h new file mode 100644 index 000000000000..237e7d90ad81 --- /dev/null +++ b/src/relay/backend/contrib/ethosn/capabilities.h @@ -0,0 +1,57 @@ +/* + * 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_RELAY_BACKEND_CONTRIB_ETHOSN_CAPABILITIES_H_ +#define TVM_RELAY_BACKEND_CONTRIB_ETHOSN_CAPABILITIES_H_ + +#include + +static std::vector targets[3] = { + { + 0x02, 0x00, 0x00, 0x00, 0x74, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x10, 0x00, 0x10, 0x00, + 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x00, + 0x10, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, + 0x01, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x10, 0x00, 0x00, + 0x00, 0x01, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x01, 0x00, + 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x40, 0x00, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x01, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + }, + { + 0x02, 0x00, 0x00, 0x00, 0x74, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x10, 0x00, 0x10, 0x00, + 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x00, + 0x10, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, + 0x01, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x10, 0x00, 0x00, + 0x00, 0x01, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x01, 0x00, + 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x40, 0x00, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x01, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + }, + { + 0x02, 0x00, 0x00, 0x00, 0x74, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x10, 0x00, 0x10, 0x00, + 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x00, + 0x10, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, + 0x01, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x10, 0x00, 0x00, + 0x00, 0x01, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x01, 0x00, + 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x40, 0x00, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x01, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + }}; + +#endif // TVM_RELAY_BACKEND_CONTRIB_ETHOSN_CAPABILITIES_H_ diff --git a/src/relay/backend/contrib/ethosn/codegen.cc b/src/relay/backend/contrib/ethosn/codegen.cc new file mode 100644 index 000000000000..b477c0ce8cfc --- /dev/null +++ b/src/relay/backend/contrib/ethosn/codegen.cc @@ -0,0 +1,214 @@ +/* + * 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 src/relay/backend/contrib/ethosn/codegen.cc + * \brief The Relay -> Ethos-N command stream compiler. + */ +#include +#include + +#include "codegen_ethosn.h" +#include "ethosn_api.h" + +namespace tvm { +namespace relay { +namespace contrib { +namespace ethosn { + +sl::TensorInfo GetTensorInfo(std::map> tensor_table, + const Call& call) { + if (tensor_table.find(call) != tensor_table.end()) return tensor_table[call][0]; + + return sl::TensorInfo(); +} + +void InferTensorsVisitor::InferCall(const CallNode* cn) { + EthosnError err; + Call call = GetRef(cn); + // Determine call -> NPU mapping + if (EthosnAPI::IsEthosOp(call, "qnn.concatenate")) { + ConcatenateParams params; + err = EthosnAPI::Concatenate(call, ¶ms); + tensor_table_[cn->args[0]] = params.input_infos; + } else if (EthosnAPI::IsEthosOp(call, "split")) { + SplitParams params; + params.input_info = GetTensorInfo(tensor_table_, call); + err = EthosnAPI::Split(call, ¶ms); + tensor_table_[cn->args[0]] = {params.input_info}; + } else { + err = EthosnError("unknown operator"); + } + if (err) { + ReportFatalError(call, err); + } +} + +// This will only visit an expression if the expression's tensor info +// has already been entirely inferred. +// An example where this is important is a tuple node where each +// get item node will only infer one field of the tuple's expression info. +// We don't want to traverse the tuple until all of its fields have been inferred. +void InferTensorsVisitor::VisitInferred(const Expr& expr) { + if (tensor_table_.find(expr) != tensor_table_.end()) { + for (const auto& tensor_info : tensor_table_[expr]) { + if (tensor_info == sl::TensorInfo()) return; + } + VisitExpr(expr); + } +} + +void InferTensorsVisitor::VisitExpr_(const CallNode* cn) { + InferCall(cn); + // Pre-order visitor + for (const auto& arg : cn->args) { + VisitInferred(arg); + } +} + +void InferTensorsVisitor::VisitExpr_(const TupleNode* tn) { + auto tuple = GetRef(tn); + CHECK(tensor_table_.find(tuple) != tensor_table_.end()); + for (size_t i = 0; i < tn->fields.size(); i++) { + tensor_table_[tn->fields[i]] = {tensor_table_[tuple][i]}; + } + // Pre-order visitor + for (const auto& field : tn->fields) { + VisitExpr(field); + } +} + +void InferTensorsVisitor::VisitExpr_(const TupleGetItemNode* tgn) { + // Don't assume it must be targeting a TupleNode + // Vars and calls can still have TupleType + auto tg = GetRef(tgn); + CHECK(tensor_table_.find(tg) != tensor_table_.end()); + auto tuple = tg->tuple; + auto type = tuple->checked_type().as(); + int index = tg->index; + // Resize the tensor infos to the tuple size if not already done + if (tensor_table_.find(tuple) == tensor_table_.end()) { + tensor_table_[tuple].resize(type->fields.size()); + } + tensor_table_[tuple][index] = tensor_table_[tg][0]; + // Pre-order visitor + VisitInferred(tuple); +} + +sl::TensorsAndId MakeOps(const sl::TensorAndId& op) { + sl::TensorsAndId ops; + ops.tensors = {op.tensor}; + ops.operationId = op.operationId; + return ops; +} + +sl::TensorsAndId ConstructNetworkVisitor::HandleCall(const CallNode* cn) { + EthosnError err; + Call call = GetRef(cn); + sl::TensorAndId tensor; + sl::TensorsAndId tensors; + // Determine call -> NPU mapping + if (EthosnAPI::IsEthosOp(call, "qnn.concatenate")) { + if ((err = MakeConcatenateLayer(call, &tensor))) ReportFatalError(call, err); + return MakeOps(tensor); + } else if (EthosnAPI::IsEthosOp(call, "split")) { + if ((err = MakeSplitLayer(call, &tensors))) ReportFatalError(call, err); + return tensors; + } else { + ReportFatalError(call, EthosnError("unknown operator")); + return {}; + } +} + +void ConstructNetworkVisitor::VisitExpr_(const CallNode* cn) { + auto operand = HandleCall(cn); + operand_table_[GetRef(cn)] = operand.tensors; + for (size_t i = 0; i < operand.tensors.size(); i++) { + id_table_[GetRef(cn)].push_back(std::make_pair(operand.operationId, i)); + } +} + +void ConstructNetworkVisitor::VisitExpr_(const TupleNode* op) { + Tuple tuple = GetRef(op); + for (const auto& arg : tuple->fields) { + // The fields in a tuple should not themselves be tuples + // Nested tuples are not supported + if (operand_table_[arg].size() == 1) { + operand_table_[tuple].push_back(operand_table_[arg][0]); + id_table_[tuple].push_back(id_table_[arg][0]); + } else { + operand_table_[tuple].push_back(nullptr); + id_table_[tuple].push_back(std::make_pair(0, 0)); + } + } +} + +void ConstructNetworkVisitor::VisitExpr_(const TupleGetItemNode* tg) { + Expr tuple = tg->tuple; + operand_table_[GetRef(tg)] = {operand_table_[tuple][tg->index]}; + id_table_[GetRef(tg)] = {id_table_[tuple][tg->index]}; +} + +void ConstructNetworkVisitor::VisitLeaf(const Expr& expr) { + // Don't traverse into functions, they're not supported + if (!expr->IsInstance()) MixedModeVisitor::VisitLeaf(expr); +} + +EthosnError ConstructNetworkVisitor::MakeConcatenateLayer(const Call& call, + sl::TensorAndId* out) { + ConcatenateParams params; + if (auto err = EthosnAPI::Concatenate(call, ¶ms)) { + return err; + } + + std::vector layers; + auto ops = operand_table_[call->args[0]]; + + for (const auto& op : ops) { + layers.emplace_back(op.get()); + } + try { + *out = AddConcatenation(network_, layers, params.concat_info); + } catch (const sl::NotSupportedException& e) { + return EthosnError(e.what()); + } + return EthosnError(); +} + +EthosnError ConstructNetworkVisitor::MakeSplitLayer(const Call& call, sl::TensorsAndId* outs) { + SplitParams params; + params.input_info = GetTensorInfo(tensor_table_, call); + if (auto err = EthosnAPI::Split(call, ¶ms)) { + return err; + } + + auto input = operand_table_[call->args[0]][0]; + + try { + *outs = AddSplit(network_, *input, params.split_info); + } catch (const sl::NotSupportedException& e) { + return EthosnError(e.what()); + } + return EthosnError(); +} + +} // namespace ethosn +} // namespace contrib +} // namespace relay +} // namespace tvm diff --git a/src/relay/backend/contrib/ethosn/codegen_ethosn.h b/src/relay/backend/contrib/ethosn/codegen_ethosn.h new file mode 100644 index 000000000000..72ed6ed86ecc --- /dev/null +++ b/src/relay/backend/contrib/ethosn/codegen_ethosn.h @@ -0,0 +1,331 @@ +/* + * 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 src/relay/backend/contrib/ethosn/codegen_ethosn.h + * \brief The Relay -> Ethos-N command stream compiler. + */ + +#ifndef TVM_RELAY_BACKEND_CONTRIB_ETHOSN_CODEGEN_ETHOSN_H_ +#define TVM_RELAY_BACKEND_CONTRIB_ETHOSN_CODEGEN_ETHOSN_H_ + +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "../../../../runtime/contrib/ethosn/ethosn_runtime.h" +#include "../codegen_c/codegen_c.h" +#include "ethosn_api.h" +#include "ethosn_support_library/Support.hpp" +#include "ethosn_support_library/SupportQueries.hpp" + +namespace tvm { +namespace relay { +namespace contrib { +namespace ethosn { + +namespace sl = ::ethosn::support_library; + +/*! + * \brief A struct to hold an uncompiled support library network alongside + * the desired order of input and output operation ids. + */ +struct NetworkWithIDs { + struct hash_pair { + template + size_t operator()(const std::pair& p) const { + return std::hash{}(p.first) ^ std::hash{}(p.second); + } + }; + std::shared_ptr network; + std::unordered_map input_ids; + std::unordered_map, unsigned int, hash_pair> output_ids; +}; + +/*! + * \brief A base class for error handling using ErrorReporter. + */ +class ErrorReportingPass { + public: + ErrorReportingPass(const IRModule& mod, const GlobalVar& var) : mod_(mod), var_(var) {} + + /*! + * \brief Report fatal errors for an expression. + * \param expr The expression to report errors at. + * \param err The errors to report. + */ + void ReportFatalError(const ObjectRef& expr, const EthosnError& err) { + for (const auto& msg : err.msgs) { + error_reporter_.ReportAt(this->var_, expr, ErrorBuilder() << msg); + } + error_reporter_.RenderErrors(this->mod_); + } + + protected: + /*! \brief An ErrorReporter object to render the errors.*/ + ErrorReporter error_reporter_; + /*! \brief The module to report errors for. */ + IRModule mod_; + /*! \brief The GlobalVar to report errors for. */ + GlobalVar var_; +}; + +/*! + * \brief A custom pass to infer the support library tensor information + * for a Relay expression. + * + * Support Library requires that tensors are explicitly declared with + * information on their size, data type, format (eg. NHWC) and quantisation + * parameters. In Relay, size and data type are already determined when the + * type_infer pass is run. However, format and quantisation parameters are + * properties of the operators that consume the tensors. + * + * This pass works by having each node initialise the information of its + * parents, essentially propagating the inferred information all the way up + * to the inputs of the expression. + * + * Because the children initialise the information of the parents, it is + * necessary to traverse the graph in such a way so as to ensure all the + * children of a node are visited before the parent is. As Relay does not + * keep a reference to child nodes, this pass goes in preorder but will + * skip visiting a parent if all the children haven't yet been visited (see + * VisitInferred for the logic that implements this). + * + * Inference only works for supported callnodes, for tuplenodes, tuplegetitem + * nodes and free var nodes. Other nodes should not be off-loaded to Ethos-N. + */ +class InferTensorsVisitor : private ErrorReportingPass, private ExprVisitor { + public: + InferTensorsVisitor(const IRModule& mod, const GlobalVar& var) : ErrorReportingPass(mod, var) {} + + /*! + * \brief Infer the support library tensor information for all the nodes + * in an expression. + * \param expr The expression for which to infer tensor information. + * \return A map of expressions to tensor information. + * \note This algorithm does not traverse into functions, so call it on + * the body of the function you're interested in. + */ + std::map> Infer(const Expr& expr) { + tensor_table_.clear(); + CHECK(expr->checked_type().defined()); + size_t output_size = 1; + if (expr->checked_type()->IsInstance()) { + auto type = expr->checked_type().as(); + output_size = type->fields.size(); + } + for (size_t i = 0; i < output_size; i++) { + tensor_table_[expr].push_back(sl::TensorInfo({1, 1, 1, 1}, sl::DataType::UINT8_QUANTIZED, + sl::DataFormat::NHWC, sl::QuantizationInfo())); + } + VisitInferred(expr); + return tensor_table_; + } + + private: + // Infer a callnode if it's a supported operator/composite function + void InferCall(const CallNode* cn); + void VisitInferred(const Expr& expr); + + void VisitExpr_(const CallNode* cn) final; + void VisitExpr_(const TupleNode* tn) final; + void VisitExpr_(const TupleGetItemNode* tg) final; + // Don't traverse into functions, the Ethos-N codegen isn't meant to support them. + void VisitExpr_(const FunctionNode* fn) final {} + + /*! \brief A look-up table from Expr to tensor infos. */ + std::map> tensor_table_; +}; + +std::map> InferTensors(const IRModule& mod, const GlobalVar& var, + const Expr& expr) { + return InferTensorsVisitor(mod, var).Infer(expr); +} + +/*! + * \brief A pass to generate a support library network from a Relay function. + * + * This pass constructs an equivalent support library network from a Relay + * function in two visits. One to infer the tensor information of all the nodes + * and another in postorder to add the nodes as support library operands. + * (Supported) Callnodes, tuplenodes, tuplegetitemnodes and (free) + * varnodes are handled by this pass. + * + * As part of the pass, nodes in the function body are associated with both + * type information in the 'tensor_table', and support library operands in the + * 'operand_table'. Both of these are maps of vectors as a Relay node can have + * tuple type and accordingly be associated with multiple tensors. For nodes + * which are not tuple type, vectors of size 1 are used. + */ +class ConstructNetworkVisitor : public MixedModeVisitor, private ErrorReportingPass { + public: + explicit ConstructNetworkVisitor(const IRModule& mod, const GlobalVar& var) + : ErrorReportingPass(mod, var) {} + + /*! + * \brief Construct a support library network from a given Relay function. The + * function should contain only nodes supported by Ethos-N. + * \param func The Relay function for which to construct a support library network. + * \return A support library network that performs the same operation as the Relay + * function. + */ + NetworkWithIDs Construct(const Function& func) { + // Initialise everything + NetworkWithIDs network_with_ids; + network_ = sl::CreateNetwork(); + network_with_ids.network = network_; + operand_table_.clear(); + + // Infer tensor information + tensor_table_ = InferTensors(this->mod_, this->var_, func->body); + // Add the inputs in the order they appear in the parameters + unsigned int idx = 0; + for (const auto& param : func->params) { + for (const auto& tensor_info : tensor_table_[param]) { + auto tensor_and_id = AddInput(network_, tensor_info); + operand_table_[param].push_back(tensor_and_id.tensor); + id_table_[param].push_back(std::make_pair(tensor_and_id.operationId, 0)); + network_with_ids.input_ids[tensor_and_id.operationId] = idx++; + } + } + // Add the function body + VisitExpr(func->body); + // Add the outputs + idx = 0; + for (const auto& layer : operand_table_[func->body]) { + AddOutput(network_, *layer); + network_with_ids.output_ids[id_table_[func->body][idx]] = idx; + idx++; + } + return network_with_ids; + } + + private: + // Translate from a callnode to the appropriate 'Make' method + sl::TensorsAndId HandleCall(const CallNode*); + + void VisitExpr_(const CallNode* cn) final; + void VisitExpr_(const TupleNode* op) final; + void VisitExpr_(const TupleGetItemNode* tg) final; + void VisitLeaf(const Expr& expr) final; + + // Make a support library operand from a Call + EthosnError MakeConcatenateLayer(const Call& call, sl::TensorAndId* out); + EthosnError MakeSplitLayer(const Call& call, sl::TensorsAndId* outs); + + /*! \brief A look-up table from Expr to layers. */ + std::map>> operand_table_; + /*! \brief A look-up table from Expr to SL operation IDs. */ + std::map>> id_table_; + /*! \brief A look-up table from Expr to tensor infos. */ + std::map> tensor_table_; + /*! \brief The support library network to compile. */ + std::shared_ptr network_; +}; + +NetworkWithIDs ConstructNetwork(const IRModule& mod, const GlobalVar& var, const Function& func) { + return ConstructNetworkVisitor(mod, var).Construct(func); +} + +class EthosnCompiler { + public: + static runtime::ethosn::OrderedCompiledNetwork CompileEthosnFunc(const IRModule& mod, + std::string name, + const Function& func) { + // Construct the network + GlobalVar var = mod->GetGlobalVar(name); + auto network_with_ids = ConstructNetwork(mod, var, func); + // Now set the required build flags + sl::CompilationOptions options = EthosnAPI::CreateOptions(); + // Finally compile the network + auto compiled_network = EthosnAPI::Compile(network_with_ids.network, options); + auto input_output_order = GetInputOutputOrder(network_with_ids, compiled_network); + runtime::ethosn::OrderedCompiledNetwork ordered_network; + ordered_network.name = name; + ordered_network.cmm = std::move(compiled_network); + ordered_network.inputs = input_output_order.first; + ordered_network.outputs = input_output_order.second; + return ordered_network; + } + + static runtime::Module CreateRuntimeModule(const ObjectRef& ref) { + std::vector cmms; + if (ref->IsInstance()) { + IRModule mod; + Function bfunc = Downcast(ref); + auto name_node = bfunc->GetAttr(tvm::attr::kGlobalSymbol); + CHECK(name_node.defined()) << "Failed to retrieved external symbol."; + mod->Add(GlobalVar(name_node.value()), bfunc); + for (const auto& it : mod->functions) { + Function func = Downcast(it.second); + name_node = func->GetAttr(tvm::attr::kGlobalSymbol); + cmms.emplace_back(CompileEthosnFunc(mod, name_node.value(), func)); + } + } else { + LOG(FATAL) << "The input ref is expected to be a Relay function or module" + << "\n"; + } + auto n = make_object(&cmms); + return runtime::Module(n); + } + + static std::pair, std::vector> GetInputOutputOrder( + NetworkWithIDs network, const std::unique_ptr& compiled_network) { + std::vector input_infos = compiled_network->GetInputBufferInfos(); + std::vector output_infos = compiled_network->GetOutputBufferInfos(); + std::vector input_order; + std::vector output_order; + for (const auto& input_info : input_infos) { + input_order.push_back(network.input_ids[input_info.m_SourceOperationId]); + } + for (const auto& output_info : output_infos) { + auto output_id = + std::make_pair(output_info.m_SourceOperationId, output_info.m_SourceOperationOutputIndex); + output_order.push_back(network.output_ids[output_id]); + } + return std::make_pair(input_order, output_order); + } +}; + +runtime::Module CompileEthosn(const ObjectRef& ref) { + return EthosnCompiler::CreateRuntimeModule(ref); +} + +TVM_REGISTER_GLOBAL("relay.ext.ethos-n").set_body_typed(CompileEthosn); + +} // namespace ethosn +} // namespace contrib +} // namespace relay +} // namespace tvm + +#endif // TVM_RELAY_BACKEND_CONTRIB_ETHOSN_CODEGEN_ETHOSN_H_ diff --git a/src/relay/backend/contrib/ethosn/ethosn_api.cc b/src/relay/backend/contrib/ethosn/ethosn_api.cc new file mode 100644 index 000000000000..df0ec1411a14 --- /dev/null +++ b/src/relay/backend/contrib/ethosn/ethosn_api.cc @@ -0,0 +1,268 @@ +/* + * 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 "ethosn_api.h" + +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include + +#include "capabilities.h" +#include "ethosn_support_library/Support.hpp" +#include "ethosn_support_library/SupportQueries.hpp" + +namespace tvm { +namespace relay { +namespace contrib { +namespace ethosn { + +std::unique_ptr EthosnAPI::Compile(std::shared_ptr network, + const sl::CompilationOptions& options) { + std::vector> compiled_network = + sl::Compile(*network, options); + CHECK_GE(compiled_network.size(), 1) << "Ethos-N compiler failed to compile network"; + + return std::move(compiled_network[0]); +} + +struct EthosnCompilerConfigNode : public tvm::AttrsNode { + int variant; + bool strategy0; + bool strategy1; + bool strategy3; + bool strategy4; + bool strategy6; + bool strategy7; + bool dump_ram; + bool initial_sram_dump; + bool block_config_16x16; + bool block_config_32x8; + bool block_config_8x32; + bool block_config_8x8; + bool enable_intermediate_compression; + bool disable_winograd; + bool dump_debug_files; + String debug_dir; + bool enable_cascading; + + TVM_DECLARE_ATTRS(EthosnCompilerConfigNode, "ext.attrs.EthosnCompilerConfigNode") { + TVM_ATTR_FIELD(variant) + .describe("0 for Ethos-N77, 1 for Ethos-N57, 2 for Ethos-N37. See Ethos-N documentation.") + .set_default(0); + TVM_ATTR_FIELD(strategy0).set_default(true); + TVM_ATTR_FIELD(strategy1).set_default(true); + TVM_ATTR_FIELD(strategy3).set_default(true); + TVM_ATTR_FIELD(strategy4).set_default(true); + TVM_ATTR_FIELD(strategy6).set_default(true); + TVM_ATTR_FIELD(strategy7).set_default(true); + TVM_ATTR_FIELD(dump_ram).set_default(false); + TVM_ATTR_FIELD(initial_sram_dump).set_default(false); + TVM_ATTR_FIELD(block_config_16x16).set_default(true); + TVM_ATTR_FIELD(block_config_32x8).set_default(true); + TVM_ATTR_FIELD(block_config_8x32).set_default(true); + TVM_ATTR_FIELD(block_config_8x8).set_default(true); + TVM_ATTR_FIELD(enable_intermediate_compression).set_default(true); + TVM_ATTR_FIELD(disable_winograd).set_default(false); + TVM_ATTR_FIELD(dump_debug_files).set_default(false); + TVM_ATTR_FIELD(debug_dir).set_default("."); + TVM_ATTR_FIELD(enable_cascading).set_default(false); + } +}; + +class EthosnCompilerConfig : public Attrs { + public: + TVM_DEFINE_NOTNULLABLE_OBJECT_REF_METHODS(EthosnCompilerConfig, Attrs, EthosnCompilerConfigNode); +}; + +TVM_REGISTER_NODE_TYPE(EthosnCompilerConfigNode); +TVM_REGISTER_PASS_CONFIG_OPTION("relay.ext.ethos-n.options", EthosnCompilerConfig); + +sl::CompilationOptions EthosnAPI::CreateOptions() { + auto ctx = transform::PassContext::Current(); + auto cfg = ctx->GetConfig("relay.ext.ethos-n.options"); + if (!cfg.defined()) { + cfg = AttrsWithDefaultValues(); + } + + sl::CompilationOptions options(targets[cfg.value()->variant]); + options.m_Strategy0 = cfg.value()->strategy0; + options.m_Strategy1 = cfg.value()->strategy1; + options.m_Strategy3 = cfg.value()->strategy3; + options.m_Strategy4 = cfg.value()->strategy4; + options.m_Strategy6 = cfg.value()->strategy6; + options.m_Strategy7 = cfg.value()->strategy7; + options.m_DebugInfo.m_DumpRam = cfg.value()->dump_ram; + options.m_DebugInfo.m_InitialSramDump = cfg.value()->initial_sram_dump; + options.m_BlockConfig16x16 = cfg.value()->block_config_16x16; + options.m_BlockConfig32x8 = cfg.value()->block_config_32x8; + options.m_BlockConfig8x32 = cfg.value()->block_config_8x32; + options.m_BlockConfig8x8 = cfg.value()->block_config_8x8; + options.m_EnableIntermediateCompression = cfg.value()->enable_intermediate_compression; + options.m_DisableWinograd = cfg.value()->disable_winograd; + options.m_DebugInfo.m_DumpDebugFiles = cfg.value()->dump_debug_files; + options.m_DebugInfo.m_DebugDir = cfg.value()->debug_dir; + options.m_EnableCascading = cfg.value()->enable_cascading; + return options; +} + +bool EthosnAPI::IsEthosFunc(const Call& call, const std::string& op_name) { + if (call->op->IsInstance()) { + Function func = Downcast(call->op); + CHECK(func.defined()); + auto name_node = func->GetAttr(attr::kComposite); + return name_node.value() == op_name; + } + return false; +} + +bool EthosnAPI::IsEthosOp(const Call& call, const std::string& op_name) { + if (call->op->IsInstance()) { + Op op = Downcast(call->op); + CHECK(op.defined()); + return op == Op::Get(op_name); + } else { + return false; + } +} + +EthosnError EthosnAPI::Concatenate(const Expr& expr, ConcatenateParams* params) { + Call call = Downcast(expr); + const auto& attrs = call->attrs.as(); + params->concat_info.m_Axis = attrs->axis; + + float output_s; + int output_zp; + EthosnError err = AsConstant(call->args[3], &output_s); + err += AsConstant(call->args[4], &output_zp); + params->concat_info.m_OutputQuantizationInfo = sl::QuantizationInfo(output_zp, output_s); + + auto input_scales = call->args[1].as()->fields; + auto input_zero_points = call->args[2].as()->fields; + auto input_tensors = call->args[0]->checked_type().as()->fields; + + int index = 0; + for (auto input_scale : input_scales) { + auto input_dtype = input_tensors[index].as(); + auto input_zero_point = input_zero_points[index]; + float scale; + int zp; + err += AsConstant(input_scale, &scale); + err += AsConstant(input_zero_point, &zp); + sl::TensorShape input_tensor_shape = {1, 1, 1, 1}; + sl::DataType input_data_type; + err += Tvm2Npu(input_dtype->shape, &input_tensor_shape); + err += Tvm2Npu(input_dtype->dtype, &input_data_type); + params->input_infos.emplace_back(sl::TensorInfo(input_tensor_shape, input_data_type, + sl::DataFormat::NHWC, + sl::QuantizationInfo(zp, scale))); + index++; + } + return err; +} + +EthosnError EthosnAPI::Split(const Expr& expr, SplitParams* params) { + Call call = Downcast(expr); + const auto* input_tensor_type = call->args[0]->checked_type().as(); + const auto& attrs = call->attrs.as(); + + sl::TensorShape input_tensor_shape = {1, 1, 1, 1}; + sl::DataType input_data_type; + EthosnError err = Tvm2Npu(input_tensor_type->shape, &input_tensor_shape); + err += Tvm2Npu(input_tensor_type->dtype, &input_data_type); + params->input_info = + sl::TensorInfo(input_tensor_shape, input_data_type, params->input_info.m_DataFormat, + params->input_info.m_QuantizationInfo); + params->split_info.m_Axis = attrs->axis; + if (attrs->indices_or_sections->IsInstance()) { + auto sections = Downcast(attrs->indices_or_sections)->value; + int size = input_tensor_shape[attrs->axis] / sections; + for (int i = 0; i < sections; i++) { + params->split_info.m_Sizes.push_back(size); + } + } else { + auto indices = Downcast>(attrs->indices_or_sections); + int last_index = 0; + for (const auto& i : indices) { + params->split_info.m_Sizes.push_back(i->value - last_index); + last_index = i->value; + } + int axis_size = input_tensor_shape[attrs->axis]; + params->split_info.m_Sizes.push_back(axis_size - last_index); + } + return err; +} + +EthosnError EthosnAPI::Tvm2Npu(const Array& shape, sl::TensorShape* npu_shape) { + EthosnError err = AsArray(shape, npu_shape); + if (npu_shape->front() != 1) { + err += EthosnError(ErrStrm() << "batch size=" << npu_shape->front() << ", batch size must = 1"); + } + return err; +} + +EthosnError EthosnAPI::Tvm2Npu(const tvm::DataType& dtype, sl::DataType* data_type) { + if (dtype.is_scalar() == 1) { + if (dtype.is_uint() && dtype.bits() == 8) { + *data_type = sl::DataType::UINT8_QUANTIZED; + return EthosnError(); + } else if (dtype.is_int() && dtype.bits() == 32) { + *data_type = sl::DataType::INT32_QUANTIZED; + return EthosnError(); + } + } + return EthosnError(ErrStrm() << "dtype=\'" << dtype << "\', dtype must be either uint8 or int32"); +} + +TVM_REGISTER_GLOBAL("relay.ethos-n.support.concatenate") + .set_body([](tvm::TVMArgs args, tvm::TVMRetValue* rv) { + Call call = args[0]; + ConcatenateParams params; + auto err = EthosnAPI::Concatenate(call, ¶ms); + *rv = !err && sl::IsConcatenationSupported(params.input_infos, params.concat_info); + }); + +TVM_REGISTER_GLOBAL("relay.ethos-n.support.split") + .set_body([](tvm::TVMArgs args, tvm::TVMRetValue* rv) { + Call call = args[0]; + SplitParams params; + auto err = EthosnAPI::Split(call, ¶ms); + *rv = !err && sl::IsSplitSupported(params.input_info, params.split_info); + }); + +TVM_REGISTER_GLOBAL("relay.ethos-n.query").set_body([](tvm::TVMArgs args, tvm::TVMRetValue* rv) { +#if defined ETHOSN_HW + *rv = true; +#else + *rv = false; +#endif +}); + +} // namespace ethosn +} // namespace contrib +} // namespace relay +} // namespace tvm diff --git a/src/relay/backend/contrib/ethosn/ethosn_api.h b/src/relay/backend/contrib/ethosn/ethosn_api.h new file mode 100644 index 000000000000..ac3179504853 --- /dev/null +++ b/src/relay/backend/contrib/ethosn/ethosn_api.h @@ -0,0 +1,142 @@ +/* + * 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_RELAY_BACKEND_CONTRIB_ETHOSN_ETHOSN_API_H_ +#define TVM_RELAY_BACKEND_CONTRIB_ETHOSN_ETHOSN_API_H_ + +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include + +#include "ethosn_support_library/Support.hpp" +#include "ethosn_support_library/SupportQueries.hpp" + +namespace tvm { +namespace relay { +namespace contrib { +namespace ethosn { + +namespace sl = ::ethosn::support_library; + +struct ConcatenateParams { + sl::QuantizationInfo qInfo; + sl::ConcatenationInfo concat_info = sl::ConcatenationInfo(1, qInfo); + std::vector input_infos; +}; + +struct SplitParams { + sl::SplitInfo split_info = sl::SplitInfo(0, {}); + sl::TensorInfo input_info; +}; + +class ErrStrm { + public: + template + ErrStrm& operator<<(const T& val) { // NOLINT(*) + stream_ << val; + return *this; + } + + private: + std::stringstream stream_; + friend class EthosnError; +}; + +class EthosnError { + public: + EthosnError() {} + explicit EthosnError(const Array& msgs) : msgs(msgs) {} + explicit EthosnError(const String& msg) { msgs.push_back(msg); } + explicit EthosnError(const ErrStrm& err) : EthosnError(err.stream_.str()) {} + + explicit operator bool() const { return !msgs.empty(); } + + EthosnError& operator+=(const EthosnError& other) { + msgs.insert(msgs.end(), other.msgs.begin(), other.msgs.end()); + return *this; + } + + Array msgs; +}; + +class EthosnAPI { + public: + static std::unique_ptr Compile(std::shared_ptr network, + const sl::CompilationOptions& options); + + static sl::CompilationOptions CreateOptions(); + + static bool IsEthosFunc(const Call& call, const std::string& op_name); + static bool IsEthosOp(const Call& call, const std::string& op_name); + + static EthosnError Concatenate(const Expr& expr, ConcatenateParams* params); + static EthosnError Split(const Expr& expr, SplitParams* params); + + private: + static EthosnError Tvm2Npu(const Array& shape, sl::TensorShape* npu_shape); + static EthosnError Tvm2Npu(const tvm::DataType& dtype, sl::DataType* data_type); + + // Convert an array of IntImmNodes into ValueT + // IndexT type of Array indexing variable + // ValueT type of resulting value + template + static EthosnError AsArray(const Array& arr, std::array* v) { + if (arr.size() > 4) + return EthosnError(ErrStrm() << "dimensions=" << arr.size() << ", dimensions must be <= 4"); + for (size_t i = 0; i < std::min(arr.size(), 4ul); i++) { + const PrimExpr& a = arr[i]; + const auto* intImm = a.as(); + if (intImm->value > std::numeric_limits::max()) { + return EthosnError(ErrStrm() << "axis size=" << intImm->value << ", axis size must be <= " + << std::numeric_limits::max()); + } + (*v)[i] = static_cast(intImm->value); + } + return EthosnError(); + } + + // Get a T from a constant represented by + // a NDArray. + template + static EthosnError AsConstant(const Expr& expr, T* out) { + if (!expr->IsInstance()) { + return EthosnError("expected constant data"); + } + runtime::NDArray data = Downcast(expr)->data; + *out = *static_cast(data.operator->()->data); + return EthosnError(); + } +}; + +} // namespace ethosn +} // namespace contrib +} // namespace relay +} // namespace tvm + +#endif // TVM_RELAY_BACKEND_CONTRIB_ETHOSN_ETHOSN_API_H_ diff --git a/src/runtime/contrib/ethosn/ethosn_device.cc b/src/runtime/contrib/ethosn/ethosn_device.cc new file mode 100644 index 000000000000..5b9e2a0e21f6 --- /dev/null +++ b/src/runtime/contrib/ethosn/ethosn_device.cc @@ -0,0 +1,222 @@ +/* + * 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 ethosn_device.cc + * \brief Ethos-N NPU device integration. + */ + +#include +#include +#include +#include + +#include +#include + +#include "ethosn_driver_library/Buffer.hpp" +#include "ethosn_support_library/Support.hpp" + +#if defined ETHOSN_HW + +#include "ethosn_driver_library/Inference.hpp" +#include "ethosn_driver_library/Network.hpp" + +namespace tvm { +namespace runtime { +namespace ethosn { + +namespace sl = ::ethosn::support_library; +namespace dl = ::ethosn::driver_library; + +int64_t GetTensorSize(const DLTensor& tensor) { + int64_t size = 1; + for (int i = 0; i < tensor.ndim; i++) { + size *= tensor.shape[i]; + } + return size; +} + +bool WaitForInference(dl::Inference* inference, int timeout) { + // Wait for inference to complete + int fd = inference->GetFileDescriptor(); + struct pollfd fds; + memset(&fds, 0, sizeof(fds)); + fds.fd = fd; + fds.events = POLLIN; // Wait for any available input. + + const int ms_per_seconds = 1000; + int poll_result = poll(&fds, 1, timeout * ms_per_seconds); + if (poll_result > 0) { + dl::InferenceResult result; + if (read(fd, &result, sizeof(result)) != sizeof(result)) { + return false; + } + if (result != dl::InferenceResult::Completed) { + return false; + } + } else if (poll_result == 0) { + return false; + } else { + return false; + } + return true; +} + +template +void CopyOutput(dl::Buffer* source_buffers[], std::vector* outputs) { + for (DLTensor* tensor : *outputs) { + dl::Buffer* source_buffer = source_buffers[0]; + uint8_t* source_buffer_data = source_buffer->GetMappedBuffer(); + size_t size = source_buffer->GetSize(); + T* dest_pointer = static_cast(tensor->data); + std::copy_backward(source_buffer_data, source_buffer_data + size, dest_pointer + size); + source_buffers++; + } +} + +void CreateBuffers(std::vector >* fm, + const std::vector& tensors) { + int index = 0; + for (auto buffer : tensors) { + auto* data = static_cast(buffer->data); + // The NPU only needs the size of the tensor * uint8_t. + auto data_size = static_cast(GetTensorSize(*buffer)); + (*fm)[index++] = std::make_shared(data, data_size, dl::DataFormat::NHWC); + } +} + +bool Inference(tvm::runtime::TVMArgs args, sl::CompiledNetwork* network, + std::vector input_order, std::vector output_order) { + // Unpack parameters + uint8_t argc = 0; + std::vector inputs(input_order.size()); + for (uint8_t i = 0; i < network->GetInputBufferInfos().size(); i++) { + inputs[input_order[i]] = args[argc++]; + } + auto out_infos = network->GetOutputBufferInfos(); + std::vector outputs(output_order.size()); + for (uint8_t i = 0; i < network->GetOutputBufferInfos().size(); i++) { + outputs[output_order[i]] = args[argc++]; + } + + // Set up input buffers + std::vector > ifm(inputs.size()); + CreateBuffers(&ifm, inputs); + + // Set up output buffers + std::vector > ofm(outputs.size()); + CreateBuffers(&ofm, outputs); + + // Raw pointers for the inference + dl::Buffer* ifm_raw[inputs.size()]; + for (size_t i = 0; i < inputs.size(); i++) { + ifm_raw[i] = ifm[i].get(); + } + dl::Buffer* ofm_raw[outputs.size()]; + for (size_t i = 0; i < outputs.size(); i++) { + ofm_raw[i] = ofm[i].get(); + } + + auto npu = std::make_unique(*network); + + // Execute the inference. + std::unique_ptr result( + npu->ScheduleInference(ifm_raw, sizeof(ifm_raw) / sizeof(ifm_raw[0]), ofm_raw, + sizeof(ofm_raw) / sizeof(ofm_raw[0]))); + bool inferenceCompleted = WaitForInference(result.get(), 60); + if (inferenceCompleted) { + switch ((outputs)[0]->dtype.bits) { + case 8: { + dl::Buffer** ofms = &ofm_raw[0]; + for (DLTensor* tensor : outputs) { + uint8_t* source_buffer_data = (*ofms++)->GetMappedBuffer(); + uint8_t* dest_pointer = static_cast(tensor->data); + if (source_buffer_data != dest_pointer) { + CopyOutput(ofm_raw, &outputs); + break; + } + } + break; + } + case 16: + CopyOutput(ofm_raw, &outputs); + break; + case 32: + CopyOutput(ofm_raw, &outputs); + break; + default: + break; + } + } + + return inferenceCompleted; +} + +} // namespace ethosn +} // namespace runtime +} // namespace tvm + +#else + +#include + +namespace tvm { +namespace runtime { +namespace ethosn { + +namespace sl = ::ethosn::support_library; + +std::vector test_outputs; + +TVM_REGISTER_GLOBAL("relay.ethos-n.test.infra.inference_result") + .set_body([](tvm::TVMArgs args, tvm::TVMRetValue* rv) { + test_outputs.clear(); + for (int argc = 1; argc < args.size(); argc++) { + const DLTensor* tensor = args[argc]; + auto shape = std::vector(tensor->shape, tensor->shape + tensor->ndim); + test_outputs.emplace_back(tvm::runtime::NDArray::Empty(shape, tensor->dtype, tensor->ctx)); + test_outputs[test_outputs.size() - 1].CopyFrom(tensor); + } + }); + +// Allow the ethos-n support code to be tested without a device +bool Inference(tvm::runtime::TVMArgs args, sl::CompiledNetwork* network, + std::vector input_order, std::vector output_order) { + std::vector outputs; + for (int argc = network->GetInputBufferInfos().size(); argc < args.size(); argc++) { + outputs.push_back(args[argc]); + } + bool rc = false; + if (test_outputs.size() == outputs.size()) { + for (auto i = 0u; i < outputs.size(); i++) { + test_outputs[i].CopyTo(outputs[i]); + } + rc = true; + } + // Clear after first usage; on-exit destructor of NDArray fails + test_outputs.clear(); + return rc; +} + +} // namespace ethosn +} // namespace runtime +} // namespace tvm + +#endif diff --git a/src/runtime/contrib/ethosn/ethosn_device.h b/src/runtime/contrib/ethosn/ethosn_device.h new file mode 100644 index 000000000000..de62f60fd384 --- /dev/null +++ b/src/runtime/contrib/ethosn/ethosn_device.h @@ -0,0 +1,44 @@ +/* + * 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 ethosn_device.h + * \brief Ethos-N NPU device integration. + */ +#ifndef TVM_RUNTIME_CONTRIB_ETHOSN_ETHOSN_DEVICE_H_ +#define TVM_RUNTIME_CONTRIB_ETHOSN_ETHOSN_DEVICE_H_ + +#include + +#include "ethosn_support_library/Support.hpp" + +namespace tvm { +namespace runtime { +namespace ethosn { + +namespace sl = ::ethosn::support_library; + +bool Inference(tvm::runtime::TVMArgs args, sl::CompiledNetwork* network, + std::vector input_order, std::vector output_order); + +} // namespace ethosn +} // namespace runtime +} // namespace tvm + +#endif // TVM_RUNTIME_CONTRIB_ETHOSN_ETHOSN_DEVICE_H_ diff --git a/src/runtime/contrib/ethosn/ethosn_runtime.cc b/src/runtime/contrib/ethosn/ethosn_runtime.cc new file mode 100644 index 000000000000..98301f235e58 --- /dev/null +++ b/src/runtime/contrib/ethosn/ethosn_runtime.cc @@ -0,0 +1,146 @@ +/* + * 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 ethosn_runtime.cc + * \brief Execution handling of Ethos-N command streams. + */ + +#include "ethosn_runtime.h" + +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include + +#include "../../file_util.h" +#include "ethosn_device.h" +#include "ethosn_driver_library/Inference.hpp" +#include "ethosn_driver_library/Network.hpp" +#include "ethosn_support_library/Support.hpp" + +namespace tvm { +namespace runtime { +namespace ethosn { + +namespace sl = ::ethosn::support_library; +namespace dl = ::ethosn::driver_library; + +EthosnModule::EthosnModule(std::vector* cmms) { + for (auto& it : *cmms) { + network_map_[it.name].name = it.name; + network_map_[it.name].cmm = std::move(it.cmm); + network_map_[it.name].inputs = it.inputs; + network_map_[it.name].outputs = it.outputs; + } +} + +PackedFunc EthosnModule::GetFunction(const std::string& name, + const ObjectPtr& sptr_to_self) { + if (network_map_.find(name) != network_map_.end()) { + return PackedFunc([sptr_to_self, this, name](TVMArgs args, TVMRetValue* rv) { + *rv = Inference(args, network_map_[name].cmm.get(), network_map_[name].inputs, + network_map_[name].outputs); + }); + } else { + return PackedFunc(); + } +} + +void EthosnModule::SaveToBinary(dmlc::Stream* stream) { + stream->Write(network_map_.size()); + for (const auto& it : network_map_) { + stream->Write(it.first); + std::stringstream ss; + it.second.cmm->Serialize(ss); + stream->Write(ss.str()); + stream->Write(it.second.inputs.size()); + stream->Write(&it.second.inputs[0], sizeof(uint32_t) * it.second.inputs.size()); + stream->Write(it.second.outputs.size()); + stream->Write(&it.second.outputs[0], sizeof(uint32_t) * it.second.outputs.size()); + } +} + +Module EthosnModule::LoadFromBinary(void* strm) { + auto stream = static_cast(strm); + size_t func_count; + // Read the number of functions + stream->Read(&func_count); + std::vector cmms; + cmms.resize(func_count); + for (unsigned int i = 0; i < func_count; i++) { + OrderedCompiledNetwork& compiled = cmms[i]; + std::string ext_symbol; + std::string cmm; + uint64_t input_size; + uint64_t output_size; + // Read the symbol name + stream->Read(&compiled.name); + // Read the serialized command stream + stream->Read(&cmm); + std::istringstream cmm_strm(cmm); + compiled.cmm = sl::DeserializeCompiledNetwork(cmm_strm); + // Read the number of inputs + stream->Read(&input_size); + auto size = static_cast(input_size); + compiled.inputs.resize(size); + // Read the order of inputs + stream->Read(&compiled.inputs[0], sizeof(uint32_t) * size); + // Read the number of outputs + stream->Read(&output_size); + size = static_cast(output_size); + compiled.outputs.resize(size); + // Read the order of outputs + stream->Read(&compiled.outputs[0], sizeof(uint32_t) * size); + } + auto n = make_object(&cmms); + return Module(n); +} + +void EthosnModule::SaveToFile(const std::string& path, const std::string& format) { + std::string data; + dmlc::MemoryStringStream writer(&data); + dmlc::SeekStream* strm = &writer; + SaveToBinary(strm); + SaveBinaryToFile(path, data); +} + +Module EthosnModule::LoadFromFile(const std::string& path) { + std::string data; + LoadBinaryFromFile(path, &data); + dmlc::MemoryStringStream reader(&data); + return LoadFromBinary(&reader); +} + +TVM_REGISTER_GLOBAL("runtime.module.loadfile_ethos-n").set_body([](TVMArgs args, TVMRetValue* rv) { + *rv = EthosnModule::LoadFromFile(args[0]); +}); + +TVM_REGISTER_GLOBAL("runtime.module.loadbinary_ethos-n") + .set_body([](TVMArgs args, TVMRetValue* rv) { *rv = EthosnModule::LoadFromBinary(args[0]); }); +} // namespace ethosn +} // namespace runtime +} // namespace tvm diff --git a/src/runtime/contrib/ethosn/ethosn_runtime.h b/src/runtime/contrib/ethosn/ethosn_runtime.h new file mode 100644 index 000000000000..6130518dbd0f --- /dev/null +++ b/src/runtime/contrib/ethosn/ethosn_runtime.h @@ -0,0 +1,110 @@ +/* + * 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 ethosn_runtime.h + * \brief Execution handling of Ethos-N command streams. + */ +#ifndef TVM_RUNTIME_CONTRIB_ETHOSN_ETHOSN_RUNTIME_H_ +#define TVM_RUNTIME_CONTRIB_ETHOSN_ETHOSN_RUNTIME_H_ + +#include + +#include +#include // NOLINT +#include // NOLINT +#include // NOLINT +#include // NOLINT +#include // NOLINT + +namespace tvm { +namespace runtime { +namespace ethosn { + +namespace sl = ::ethosn::support_library; + +struct OrderedCompiledNetwork { + std::unique_ptr cmm; + std::string name; + std::vector inputs; + std::vector outputs; +}; + +class EthosnModule : public ModuleNode { + public: + /*! + * \brief The Ethos-N runtime module. + * \param cmms A vector of compiled networks with input/output orders. + */ + explicit EthosnModule(std::vector* cmms); + + /*! + * \brief Get a PackedFunc from the Ethos-N module. + * \param name The name of the function. + * \param sptr_to_self The ObjectPtr that points to this module node. + * \return The function pointer when it is found, otherwise, PackedFunc(nullptr). + */ + PackedFunc GetFunction(const std::string& name, const ObjectPtr& sptr_to_self) final; + /*! + * \brief Save a compiled network to a binary stream, which can then be + * serialized to disk. + * \param stream The stream to save the binary. + * \note See EthosnModule::LoadFromBinary for the serialization format. + */ + void SaveToBinary(dmlc::Stream* stream) final; + /*! + * \brief Load a compiled network from stream. + * \param strm The binary stream to load. + * \return The created Ethos-N module. + * \note The serialization format is: + * + * size_t : number of functions + * [ + * std::string : name of function (symbol) + * std::string : serialized command stream + * size_t : number of inputs + * std::vector : order of inputs + * size_t : number of outputs + * std::vector : order of outputs + * ] * number of functions + */ + static Module LoadFromBinary(void* strm); + /*! + * \brief Save a module to a specified path. + * \param path Where to save the serialized module. + */ + void SaveToFile(const std::string& path, const std::string& format) override; + /*! + * \brief Create a module from a file. + * \param path The path of the file containing the serialized module. + * \return The created Ethos-N module. + */ + static Module LoadFromFile(const std::string& path); + + const char* type_key() const override { return "ethos-n"; } + + private: + /*! \brief A map between ext_symbols (function names) and ordered compiled networks. */ + std::map network_map_; +}; + +} // namespace ethosn +} // namespace runtime +} // namespace tvm +#endif // TVM_RUNTIME_CONTRIB_ETHOSN_ETHOSN_RUNTIME_H_ diff --git a/tests/python/contrib/test_ethosn/__init__.py b/tests/python/contrib/test_ethosn/__init__.py new file mode 100644 index 000000000000..deba5e5eb494 --- /dev/null +++ b/tests/python/contrib/test_ethosn/__init__.py @@ -0,0 +1,18 @@ +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you under the Apache License, Version 2.0 (the +# "License"); you may not use this file except in compliance +# with the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, +# software distributed under the License is distributed on an +# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +# KIND, either express or implied. See the License for the +# specific language governing permissions and limitations +# under the License. +"""Infrastructure and tests for EthosN""" + diff --git a/tests/python/contrib/test_ethosn/_infrastructure.py b/tests/python/contrib/test_ethosn/_infrastructure.py new file mode 100644 index 000000000000..a71ab3dbc663 --- /dev/null +++ b/tests/python/contrib/test_ethosn/_infrastructure.py @@ -0,0 +1,22 @@ +# 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. + +"""Expose test functions to the Python front end""" + +import tvm._ffi + +tvm._ffi._init_api("relay.ethos-n.test.infra", __name__) diff --git a/tests/python/contrib/test_ethosn/infrastructure.py b/tests/python/contrib/test_ethosn/infrastructure.py new file mode 100644 index 000000000000..7345f044762c --- /dev/null +++ b/tests/python/contrib/test_ethosn/infrastructure.py @@ -0,0 +1,225 @@ +# 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. + +"""Expose Ethos test functions to the Python front end""" + +from __future__ import absolute_import, print_function +import tvm +from tvm import relay +from tvm.contrib import util, graph_runtime, download +from tvm.relay.testing import run_opt_pass +from enum import Enum +from hashlib import md5 +from itertools import zip_longest, combinations +import numpy as np +from PIL import Image +import os + +from . import _infrastructure +from tvm.relay.op.contrib import get_pattern_table + + +class Available(Enum): + UNAVAILABLE = 0 + SW_ONLY = 1 + SW_AND_HW = 2 + + +def ethosn_available(): + """Return whether Ethos-N software and hardware support is available""" + if not tvm.get_global_func("relay.ethos-n.query", True): + print("skip because Ethos-N module is not available") + return Available.UNAVAILABLE + else: + hw = tvm.get_global_func("relay.ethos-n.query")() + return Available.SW_AND_HW if hw else Available.SW_ONLY + + +def get_real_image(im_height, im_width): + repo_base = 'https://github.com/dmlc/web-data/raw/master/tensorflow/models/InceptionV1/' + img_name = 'elephant-299.jpg' + image_url = os.path.join(repo_base, img_name) + img_path = download.download_testdata(image_url, img_name, module='data') + image = Image.open(img_path).resize((im_height, im_width)) + x = np.array(image).astype('uint8') + data = np.reshape(x, (1, im_height, im_width, 3)) + return data + + +def assert_lib_hash(lib, golden): + temp = util.tempdir() + path = temp.relpath("lib.cmm") + lib.imported_modules[1].save(path) + lib_hash = md5(open(path, 'rb').read()).hexdigest() + assert lib_hash == golden, "Expected hash: {} Got hash: {}".format(golden, lib_hash) + + +def make_module(func, params): + func = relay.Function(relay.analysis.free_vars(func), func) + if len(params): + relay.build_module.bind_params_by_name(func, params) + return tvm.IRModule.from_expr(func) + + +def make_ethosn_composite(ethosn_expr, name): + vars = relay.analysis.free_vars(ethosn_expr) + func = relay.Function([relay.Var("a")], ethosn_expr) + func = func.with_attr("Composite", name) + call = relay.Call(func, vars) + return call + + +def make_ethosn_partition(ethosn_expr): + # Create an Ethos-N global function + mod = tvm.IRModule({}) + vars = relay.analysis.free_vars(ethosn_expr) + func = relay.Function(vars, ethosn_expr) + func = func.with_attr("Primitive", tvm.tir.IntImm("int32", 1)) + func = func.with_attr("Inline", tvm.tir.IntImm("int32", 1)) + func = func.with_attr("Compiler", "ethos-n") + func = func.with_attr("global_symbol", "ethos-n_0") + g1 = relay.GlobalVar("ethos-n_0") + mod[g1] = func + + # These are the vars to call the Ethos-N partition with + more_vars = relay.analysis.free_vars(ethosn_expr) + # Call the Ethos-N partition in main + call_fn1 = g1(*more_vars) + mod["main"] = relay.Function(more_vars, call_fn1) + return mod + + +def get_cpu_op_count(mod): + class Counter(tvm.relay.ExprVisitor): + def __init__(self): + super().__init__() + self.count = 0 + + def visit_call(self, call): + if isinstance(call.op, tvm.ir.Op): + self.count += 1 + + super().visit_call(call) + + c = Counter() + c.visit(mod["main"]) + return c.count + + +def build(mod, params, npu=True, cpu_ops=0, npu_partitions=1): + relay.backend.compile_engine.get().clear() + with tvm.transform.PassContext(opt_level=3, config={ + "relay.ext.ethos-n.options": {"variant": 0} + }): + with tvm.target.create("llvm -mcpu=core-avx2"): + if npu: + f = relay.build_module.bind_params_by_name(mod["main"], params) + mod = tvm.IRModule() + mod["main"] = f + mod = relay.transform.AnnotateTarget("ethos-n")(mod) + mod = relay.transform.MergeCompilerRegions()(mod) + mod = relay.transform.PartitionGraph()(mod) + cpu_op_count = get_cpu_op_count(mod) + assert cpu_op_count == cpu_ops, \ + "Got {} CPU operators, expected {}".format(cpu_op_count, cpu_ops) + partition_count = 0 + for global_var in mod.get_global_vars(): + if "ethos-n" in global_var.name_hint: + partition_count += 1 + + assert npu_partitions == partition_count, \ + "Got {} ethos-n partitions, expected {}".format(partition_count, npu_partitions) + + return relay.build(mod, params=params) + + +def run(graph, lib, params, inputs, outputs, npu=True): + module = graph_runtime.create(graph, lib, tvm.cpu()) + module.set_input(**inputs) + module.set_input(**params) + module.run() + out = [module.get_output(i) for i in range(outputs)] + if not npu: + inference_result(0, out) + return out + + +def build_and_run(mod, inputs, outputs, params, ctx=tvm.cpu(), npu=True, cpu_ops=0, npu_partitions=1): + graph, lib, params = build(mod, params, npu, cpu_ops, npu_partitions) + return run(graph, lib, params, inputs, outputs, npu) + + +def verify(answers, atol, rtol=1e-07, verify_saturation=True): + """Compare the array of answers. Each entry is a list of outputs""" + if len(answers) < 2: + print("No results to compare: expected at least two, found ", + len(answers)) + for answer in zip_longest(*answers): + for outs in combinations(answer, 2): + if verify_saturation: + assert np.count_nonzero(outs[0].asnumpy() == 255) < 0.25 * outs[0].asnumpy().size, \ + "Output is saturated: {}".format(outs[0]) + assert np.count_nonzero(outs[0].asnumpy() == 0) < 0.25 * outs[0].asnumpy().size, \ + "Output is saturated: {}".format(outs[0]) + tvm.testing.assert_allclose( + outs[0].asnumpy(), outs[1].asnumpy(), rtol=rtol, atol=atol + ) + + +def inference_result(checksum, outputs): + """Set the expected results of an Ethos inference, if the testing + infrastructure is available. This assumes that the entire graph + was offloaded to the neural processor.""" + if tvm.get_global_func( + "relay.ethos-n.test.infra.inference_result", True): + return _infrastructure.inference_result(checksum, *outputs) + return False + + +def generate_trials(space, r_factor=3): + np.random.seed(0) + max_len = 1 + for option in space: + max_len = max(max_len, len(option)) + + num_trials = r_factor * max_len + trials = [] + for i in range(num_trials): + trial = [] + for option in space: + if i % len(option) == 0: + np.random.shuffle(option) + trial.append(option[i % len(option)]) + + trials.append(trial) + + return trials + + +def test_error(mod, params, err_msg): + caught = None + with tvm.transform.PassContext(opt_level=3): + with tvm.target.create("llvm"): + try: + relay.build(mod, params) + except tvm.error.TVMError as e: + caught = e.args[0] + finally: + relay.backend.compile_engine.get().clear() + + assert caught is not None + assert err_msg in caught, caught diff --git a/tests/python/contrib/test_ethosn/test_concatenate.py b/tests/python/contrib/test_ethosn/test_concatenate.py new file mode 100644 index 000000000000..2b6e0d9bf086 --- /dev/null +++ b/tests/python/contrib/test_ethosn/test_concatenate.py @@ -0,0 +1,90 @@ +# 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. + +"""Concatenate tests for Ethos-N""" + +import numpy as np +import tvm +from tvm import relay +from . import infrastructure as tei + + +def _get_inputs(shapes): + inputs = {} + for i, shape in enumerate(shapes): + inputs["in" + str(i)] = tvm.nd.array( + np.random.randint(0, high=256, size=shape, dtype="uint8") + ) + + return inputs + + +def _get_model(shapes, dtype, axis): + tup = [] + for i, shape in enumerate(shapes): + a = relay.var("in" + str(i), shape=shape, dtype=dtype) + tup.append(a) + + zeroi = relay.const(1, "int32") + zerof = relay.const(0.5, "float32") + con = relay.qnn.op.concatenate(tup, + input_scales=[zerof]*len(shapes), + input_zero_points=[zeroi]*len(shapes), + output_scale=zerof, + output_zero_point=zeroi, + axis=axis) + return con + + +def test_concatenate(): + if not tei.ethosn_available(): + return + + trials = [ + ([(1, 4), (1, 6)], 1), + ([(1, 16, 4), (1, 16, 4)], 1), + ([(1, 25, 4, 16)]*3, 3), + ([(1, 25, 4, 16), (1, 25, 5, 16), (1, 25, 6, 16)], 2), + ] + + for shapes, axis in trials: + outputs = [] + inputs = _get_inputs(shapes) + for npu in [False, True]: + model = _get_model(shapes, "uint8", axis) + mod = tei.make_module(model, {}) + outputs.append(tei.build_and_run(mod, inputs, 1, {}, npu=npu)) + + tei.verify(outputs, 0) + + +def test_concatenate_failure(): + if not tei.ethosn_available(): + return + + trials = [ + ([(1, 4, 4, 4, 4), (1, 4, 4, 4, 4)], "uint8", 1, "dimensions=5, dimensions must be <= 4;"), + ([(1, 4, 4, 4), (1, 4, 4, 4)], "uint8", 3, "Concatenation along the channels dimension (axis 3) requires input tensors with a multiple of 16 channels;"), + ([(1, 4, 4, 4), (1, 4, 4, 4)], "int8", 2, "dtype='int8', dtype must be either uint8 or int32; dtype='int8', dtype must be either uint8 or int32;"), + ([(2, 4, 4, 4), (2, 4, 4, 4)], "uint8", 2, "batch size=2, batch size must = 1; batch size=2, batch size must = 1;"), + ([(1, 4, 4, 4), (1, 4, 4, 4)], "uint8", 0, "Concatenation cannot be performed along batch axis (axis 0);"), + ] + + for shapes, dtype, axis, err_msg in trials: + model = _get_model(shapes, dtype, axis) + mod = tei.make_ethosn_partition(model) + tei.test_error(mod, {}, err_msg) diff --git a/tests/python/contrib/test_ethosn/test_split.py b/tests/python/contrib/test_ethosn/test_split.py new file mode 100644 index 000000000000..0653dd306a30 --- /dev/null +++ b/tests/python/contrib/test_ethosn/test_split.py @@ -0,0 +1,70 @@ +# 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. + +"""Split tests for Ethos-N""" + +import numpy as np +import tvm +from tvm import relay +from . import infrastructure as tei + + +def _get_model(shape, dtype, splits, axis): + a = relay.var("a", shape=shape, dtype=dtype) + split = relay.op.split(a, indices_or_sections=splits, axis=axis) + return split.astuple() + + +def test_split(): + if not tei.ethosn_available(): + return + + trials = [ + ((1, 16, 16, 32), (2, 7, 10), 2), + ((1, 12, 8, 16), 3, 1), + ((1, 33), 11, 1), + ] + + np.random.seed(0) + for shape, splits, axis in trials: + outputs = [] + inputs = {"a": tvm.nd.array(np.random.randint(0, high=256, size=shape, dtype="uint8"))} + for npu in [False, True]: + model = _get_model(shape, "uint8", splits, axis) + mod = tei.make_module(model, {}) + output_count = splits if type(splits) == int else len(splits) + 1 + outputs.append(tei.build_and_run(mod, inputs, output_count, {}, npu=npu)) + + tei.verify(outputs, 0) + + +def test_split_failure(): + if not tei.ethosn_available(): + return + + trials = [ + ((1, 4, 4, 4, 4), "uint8", 4, 2, "dimensions=5, dimensions must be <= 4;"), + ((1, 4, 4, 4), "int8", 4, 2, "dtype='int8', dtype must be either uint8 or int32;"), + ((2, 4, 4, 4), "uint8", 4, 2, "batch size=2, batch size must = 1;"), + ((1, 4, 4, 4), "uint8", 1, 0, "Split cannot be performed along batch axis (axis 0);"), + ((1, 4, 4, 4), "uint8", 4, 3, "Split along the channels dimension (axis 3) requires all output sizes (specified in splitInfo.m_Sizes) to be multiples of 16;"), + ] + + for shape, dtype, splits, axis, err_msg in trials: + model = _get_model(shape, dtype, splits, axis) + mod = tei.make_ethosn_partition(model) + tei.test_error(mod, {}, err_msg) diff --git a/tests/python/contrib/test_ethosn/test_topologies.py b/tests/python/contrib/test_ethosn/test_topologies.py new file mode 100644 index 000000000000..435e7841a0b3 --- /dev/null +++ b/tests/python/contrib/test_ethosn/test_topologies.py @@ -0,0 +1,122 @@ +# 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. +"""Ethos-N tests for complex network topologies.""" + +import numpy as np +import tvm +from tvm import relay +from . import infrastructure as tei + + +def test_split_with_asym_concats(): + if not tei.ethosn_available(): + return + + def get_model(shape, splits, axis): + a = relay.var("a", shape=shape, dtype="uint8") + split = relay.op.split(a, indices_or_sections=splits, axis=axis) + zeroi = relay.const(1, "int32") + zerof = relay.const(0.5, "float32") + con1 = relay.qnn.op.concatenate([split[0], split[1]], + input_scales=[zerof]*2, + input_zero_points=[zeroi]*2, + output_scale=zerof, + output_zero_point=zeroi, + axis=axis) + con2 = relay.qnn.op.concatenate([split[2], split[3]], + input_scales=[zerof]*2, + input_zero_points=[zeroi]*2, + output_scale=zerof, + output_zero_point=zeroi, + axis=axis) + return relay.Tuple((con2, con1)) + + trials = [ + ((1, 16, 16, 32), (2, 7, 10), 2), + ] + + np.random.seed(0) + for shape, splits, axis in trials: + outputs = [] + inputs = {"a": tvm.nd.array(np.random.randint(0, high=256, size=shape, dtype="uint8"))} + for npu in [False, True]: + model = get_model(shape, splits, axis) + mod = tei.make_module(model, {}) + outputs.append(tei.build_and_run(mod, inputs, 2, {}, npu=npu)) + + tei.verify(outputs, 0) + + +def test_output_tuple_propagation(): + """This tests the case where the output tuple must be inferred + as having dummy tensor information.""" + if not tei.ethosn_available(): + return + + def get_model(): + a = relay.var("a", shape=(1, 4, 4, 16), dtype="uint8") + split = relay.op.split(a, indices_or_sections=4, axis=2) + return relay.Tuple((split[0], split[1], split[2], split[3])) + + np.random.seed(0) + outputs = [] + inputs = {"a": tvm.nd.array(np.random.randint(0, high=256, size=(1, 4, 4, 16), dtype="uint8"))} + for npu in [False, True]: + model = get_model() + mod = tei.make_module(model, {}) + outputs.append(tei.build_and_run(mod, inputs, 4, {}, npu=npu)) + + tei.verify(outputs, 0) + + +def test_input_tuples(): + if not tei.ethosn_available(): + return + + def get_model(shapes, axis): + tup = [] + for i, shape in enumerate(shapes): + a = relay.var("in" + str(i), shape=shape, dtype="uint8") + tup.append(a) + + zeroi = relay.const(1, "int32") + zerof = relay.const(0.5, "float32") + con = relay.qnn.op.concatenate(tup, + input_scales=[zerof]*len(shapes), + input_zero_points=[zeroi]*len(shapes), + output_scale=zerof, + output_zero_point=zeroi, + axis=axis) + + return con + + np.random.seed(0) + inputs = { + "in0": tvm.nd.array(np.random.randint(0, high=256, size=(1, 4), dtype="uint8")), + "in1": tvm.nd.array(np.random.randint(0, high=256, size=(1, 6), dtype="uint8")), + } + outputs = [] + for npu in [False, True]: + model = get_model([(1, 4), (1, 6)], 1) + if not npu: + mod = tei.make_module(model, {}) + else: + mod = tei.make_ethosn_partition(model) + graph, lib, params = tei.build(mod, {}, npu=False) + outputs.append(tei.run(graph, lib, {}, inputs, 1, npu=npu)) + + tei.verify(outputs, 0) diff --git a/tests/scripts/task_config_build_cpu.sh b/tests/scripts/task_config_build_cpu.sh index f36c1d974b7e..89e6695b83f7 100755 --- a/tests/scripts/task_config_build_cpu.sh +++ b/tests/scripts/task_config_build_cpu.sh @@ -43,3 +43,5 @@ echo set\(USE_VTA_FSIM ON\) >> config.cmake echo set\(USE_TFLITE ON\) >> config.cmake echo set\(USE_TENSORFLOW_PATH \"/tensorflow\"\) >> config.cmake echo set\(USE_FLATBUFFERS_PATH \"/flatbuffers\"\) >> config.cmake +echo set\(USE_ETHOSN /opt/arm/ethosn-driver-dev/ethos-n77\) >> config.cmake +echo set\(USE_ETHOSN_HW ON\) >> config.cmake \ No newline at end of file diff --git a/tests/scripts/task_python_ethosn_tests.sh b/tests/scripts/task_python_ethosn_tests.sh new file mode 100755 index 000000000000..36a3d0919650 --- /dev/null +++ b/tests/scripts/task_python_ethosn_tests.sh @@ -0,0 +1,30 @@ +#!/bin/bash +# 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. + +set -e +set -u +source tests/scripts/setup-pytest-env.sh + + +# Rebuild cython + +find . -type f -path "*.pyc" | xargs rm -f +make cython3 + +TVM_FFI=ctypes python3 -m pytest tests/python/contrib/test_ethosn + From a73e52412b560d400c6d0744105f23e41d3e5594 Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Fri, 7 Aug 2020 14:56:48 +0100 Subject: [PATCH 02/37] Turn off USE_ETHOSN_HW by default Change-Id: Ie2ce4528e16e93aa83df46f8a229c0ce89b45252 --- cmake/config.cmake | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cmake/config.cmake b/cmake/config.cmake index 3bd074d22851..f4a77f577154 100644 --- a/cmake/config.cmake +++ b/cmake/config.cmake @@ -218,7 +218,7 @@ set(USE_ARM_COMPUTE_LIB_GRAPH_RUNTIME OFF) set(USE_ETHOSN OFF) # If USE_ETHOSN is enabled, use Ethos-N hardware (ON) or # software test infrastructure (OFF) -set(USE_ETHOSN_HW ON) +set(USE_ETHOSN_HW OFF) # Build ANTLR parser for Relay text format # Possible values: From e1168239f949d3010eceb5f5e6075edca9ffb67c Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Fri, 7 Aug 2020 15:13:45 +0100 Subject: [PATCH 03/37] Update capabilities file Change-Id: Iebd0c62d6bc7e446662abdee4882ac874ad98aa3 --- src/relay/backend/contrib/ethosn/capabilities.h | 12 +++++++++++- src/relay/backend/contrib/ethosn/ethosn_api.cc | 2 +- 2 files changed, 12 insertions(+), 2 deletions(-) diff --git a/src/relay/backend/contrib/ethosn/capabilities.h b/src/relay/backend/contrib/ethosn/capabilities.h index 237e7d90ad81..71eb13be4691 100644 --- a/src/relay/backend/contrib/ethosn/capabilities.h +++ b/src/relay/backend/contrib/ethosn/capabilities.h @@ -22,7 +22,12 @@ #include -static std::vector targets[3] = { +namespace tvm { +namespace relay { +namespace contrib { +namespace ethosn { + +static std::vector variants[3] = { { 0x02, 0x00, 0x00, 0x00, 0x74, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x10, 0x00, 0x10, 0x00, @@ -54,4 +59,9 @@ static std::vector targets[3] = { 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, }}; +} // namespace ethosn +} // namespace contrib +} // namespace relay +} // namespace tvm + #endif // TVM_RELAY_BACKEND_CONTRIB_ETHOSN_CAPABILITIES_H_ diff --git a/src/relay/backend/contrib/ethosn/ethosn_api.cc b/src/relay/backend/contrib/ethosn/ethosn_api.cc index df0ec1411a14..755ae6765fe8 100644 --- a/src/relay/backend/contrib/ethosn/ethosn_api.cc +++ b/src/relay/backend/contrib/ethosn/ethosn_api.cc @@ -109,7 +109,7 @@ sl::CompilationOptions EthosnAPI::CreateOptions() { cfg = AttrsWithDefaultValues(); } - sl::CompilationOptions options(targets[cfg.value()->variant]); + sl::CompilationOptions options(variants[cfg.value()->variant]); options.m_Strategy0 = cfg.value()->strategy0; options.m_Strategy1 = cfg.value()->strategy1; options.m_Strategy3 = cfg.value()->strategy3; From 5350901a0c75a318dafa8d5633b42f80bdbb2b07 Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Mon, 10 Aug 2020 10:58:53 +0100 Subject: [PATCH 04/37] Fix missing header Change-Id: I0c89e380dd1d795755a1884c06a7b317a99fe297 --- src/runtime/contrib/ethosn/ethosn_device.cc | 1 + 1 file changed, 1 insertion(+) diff --git a/src/runtime/contrib/ethosn/ethosn_device.cc b/src/runtime/contrib/ethosn/ethosn_device.cc index 5b9e2a0e21f6..69139956c963 100644 --- a/src/runtime/contrib/ethosn/ethosn_device.cc +++ b/src/runtime/contrib/ethosn/ethosn_device.cc @@ -176,6 +176,7 @@ bool Inference(tvm::runtime::TVMArgs args, sl::CompiledNetwork* network, #else #include +#include namespace tvm { namespace runtime { From b69b7a853a1702d72946cb0556975e09b5e4ee33 Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Mon, 10 Aug 2020 10:59:05 +0100 Subject: [PATCH 05/37] Update cmake comments on ETHOSN_HW Change-Id: I2e96a1c818a82e5174fd94e483b0bdb3e4375a7d --- cmake/config.cmake | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cmake/config.cmake b/cmake/config.cmake index f4a77f577154..e08236c32cae 100644 --- a/cmake/config.cmake +++ b/cmake/config.cmake @@ -216,8 +216,8 @@ set(USE_ARM_COMPUTE_LIB_GRAPH_RUNTIME OFF) # - path/to/arm-ethos-N-stack: use a specific version of the # Ethos-N driver stack set(USE_ETHOSN OFF) -# If USE_ETHOSN is enabled, use Ethos-N hardware (ON) or -# software test infrastructure (OFF) +# If USE_ETHOSN is enabled, use ETHOSN_HW (ON) if Ethos-N hardware is available on this machine +# otherwise use ETHOSN_HW (OFF) to use the software test infrastructure set(USE_ETHOSN_HW OFF) # Build ANTLR parser for Relay text format From 1d03ac903ae28d64864a045387f7cede7e0a31df Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Mon, 10 Aug 2020 11:22:41 +0100 Subject: [PATCH 06/37] Add checker for case when USE_ETHOSN=OFF and USE_ETHOSN_HW=ON Change-Id: Id5c9cfb866914a0298b44ead40fcbe3764ce443c --- cmake/modules/contrib/EthosN.cmake | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/cmake/modules/contrib/EthosN.cmake b/cmake/modules/contrib/EthosN.cmake index 685ea22bebcf..7c694822b21e 100644 --- a/cmake/modules/contrib/EthosN.cmake +++ b/cmake/modules/contrib/EthosN.cmake @@ -51,4 +51,8 @@ if(NOT USE_ETHOSN STREQUAL "OFF") PROPERTIES COMPILE_FLAGS "-fno-rtti") endif() endif(ETHOSN_FOUND) +else() + if(USE_ETHOSN_HW) + message(FATAL_ERROR "Cannot enable Ethos-N HW if USE_ETHOSN=OFF") + endif() endif(NOT USE_ETHOSN STREQUAL "OFF") From 56b3294b5f48b65d1d0c64e545eeadbb4bf3334d Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Mon, 10 Aug 2020 12:01:29 +0100 Subject: [PATCH 07/37] Fix 'available' boolean Change-Id: I78e54fb9f472d2815886bea4d94b7247e0d129de --- tests/python/contrib/test_ethosn/infrastructure.py | 3 +++ 1 file changed, 3 insertions(+) diff --git a/tests/python/contrib/test_ethosn/infrastructure.py b/tests/python/contrib/test_ethosn/infrastructure.py index 7345f044762c..ed23ca9d2985 100644 --- a/tests/python/contrib/test_ethosn/infrastructure.py +++ b/tests/python/contrib/test_ethosn/infrastructure.py @@ -38,6 +38,9 @@ class Available(Enum): SW_ONLY = 1 SW_AND_HW = 2 + def __bool__(self): + return self != Available.UNAVAILABLE + def ethosn_available(): """Return whether Ethos-N software and hardware support is available""" From 73d70e6c6ac86106a9b6d833ebfb40c8789520cc Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Mon, 10 Aug 2020 12:17:53 +0100 Subject: [PATCH 08/37] Check availability in op registration Change-Id: Iecfea7dca7301dd684199c9b32f99f2113fdfd56 --- python/tvm/relay/op/contrib/ethosn.py | 26 +++++++++++++++++++ .../contrib/test_ethosn/infrastructure.py | 19 -------------- .../contrib/test_ethosn/test_concatenate.py | 5 ++-- .../python/contrib/test_ethosn/test_split.py | 5 ++-- .../contrib/test_ethosn/test_topologies.py | 7 ++--- 5 files changed, 36 insertions(+), 26 deletions(-) diff --git a/python/tvm/relay/op/contrib/ethosn.py b/python/tvm/relay/op/contrib/ethosn.py index 3f0121cd2b5e..1b14bdeb6a47 100644 --- a/python/tvm/relay/op/contrib/ethosn.py +++ b/python/tvm/relay/op/contrib/ethosn.py @@ -17,13 +17,36 @@ # pylint: disable=invalid-name, unused-argument """Arm(R) Ethos(TM) -N NPU supported operators.""" import tvm.ir +from enum import Enum from ... import qnn as _qnn from . import _ethosn as support +class Available(Enum): + UNAVAILABLE = 0 + SW_ONLY = 1 + SW_AND_HW = 2 + + def __bool__(self): + return self != Available.UNAVAILABLE + + +def ethosn_available(): + """Return whether Ethos-N software and hardware support is available""" + if not tvm.get_global_func("relay.ethos-n.query", True): + print("skip because Ethos-N module is not available") + return Available.UNAVAILABLE + else: + hw = tvm.get_global_func("relay.ethos-n.query")() + return Available.SW_AND_HW if hw else Available.SW_ONLY + + @tvm.ir.register_op_attr("qnn.concatenate", "target.ethos-n") def qnn_concatenate(attrs, args): """Check if a concatenate is supported by Ethos-N.""" + if not ethosn_available(): + return False + conc = _qnn.op.concatenate(*args, **attrs) if not support.concatenate(conc): return False @@ -50,6 +73,9 @@ def qnn_concatenate(attrs, args): @tvm.ir.register_op_attr("split", "target.ethos-n") def split(attrs, args): """Check if a split is supported by Ethos-N.""" + if not ethosn_available(): + return False + if isinstance(attrs["indices_or_sections"], tvm.tir.IntImm): sp = tvm.relay.split(*args, indices_or_sections=attrs["indices_or_sections"].value, diff --git a/tests/python/contrib/test_ethosn/infrastructure.py b/tests/python/contrib/test_ethosn/infrastructure.py index ed23ca9d2985..9dd33345c370 100644 --- a/tests/python/contrib/test_ethosn/infrastructure.py +++ b/tests/python/contrib/test_ethosn/infrastructure.py @@ -33,25 +33,6 @@ from tvm.relay.op.contrib import get_pattern_table -class Available(Enum): - UNAVAILABLE = 0 - SW_ONLY = 1 - SW_AND_HW = 2 - - def __bool__(self): - return self != Available.UNAVAILABLE - - -def ethosn_available(): - """Return whether Ethos-N software and hardware support is available""" - if not tvm.get_global_func("relay.ethos-n.query", True): - print("skip because Ethos-N module is not available") - return Available.UNAVAILABLE - else: - hw = tvm.get_global_func("relay.ethos-n.query")() - return Available.SW_AND_HW if hw else Available.SW_ONLY - - def get_real_image(im_height, im_width): repo_base = 'https://github.com/dmlc/web-data/raw/master/tensorflow/models/InceptionV1/' img_name = 'elephant-299.jpg' diff --git a/tests/python/contrib/test_ethosn/test_concatenate.py b/tests/python/contrib/test_ethosn/test_concatenate.py index 2b6e0d9bf086..cca61d1db677 100644 --- a/tests/python/contrib/test_ethosn/test_concatenate.py +++ b/tests/python/contrib/test_ethosn/test_concatenate.py @@ -20,6 +20,7 @@ import numpy as np import tvm from tvm import relay +from tvm.relay.op.contrib.ethosn import ethosn_available from . import infrastructure as tei @@ -51,7 +52,7 @@ def _get_model(shapes, dtype, axis): def test_concatenate(): - if not tei.ethosn_available(): + if not ethosn_available(): return trials = [ @@ -73,7 +74,7 @@ def test_concatenate(): def test_concatenate_failure(): - if not tei.ethosn_available(): + if not ethosn_available(): return trials = [ diff --git a/tests/python/contrib/test_ethosn/test_split.py b/tests/python/contrib/test_ethosn/test_split.py index 0653dd306a30..d5ff9bf0831f 100644 --- a/tests/python/contrib/test_ethosn/test_split.py +++ b/tests/python/contrib/test_ethosn/test_split.py @@ -20,6 +20,7 @@ import numpy as np import tvm from tvm import relay +from tvm.relay.op.contrib.ethosn import ethosn_available from . import infrastructure as tei @@ -30,7 +31,7 @@ def _get_model(shape, dtype, splits, axis): def test_split(): - if not tei.ethosn_available(): + if not ethosn_available(): return trials = [ @@ -53,7 +54,7 @@ def test_split(): def test_split_failure(): - if not tei.ethosn_available(): + if not ethosn_available(): return trials = [ diff --git a/tests/python/contrib/test_ethosn/test_topologies.py b/tests/python/contrib/test_ethosn/test_topologies.py index 435e7841a0b3..942186d57e60 100644 --- a/tests/python/contrib/test_ethosn/test_topologies.py +++ b/tests/python/contrib/test_ethosn/test_topologies.py @@ -19,11 +19,12 @@ import numpy as np import tvm from tvm import relay +from tvm.relay.op.contrib.ethosn import ethosn_available from . import infrastructure as tei def test_split_with_asym_concats(): - if not tei.ethosn_available(): + if not ethosn_available(): return def get_model(shape, splits, axis): @@ -64,7 +65,7 @@ def get_model(shape, splits, axis): def test_output_tuple_propagation(): """This tests the case where the output tuple must be inferred as having dummy tensor information.""" - if not tei.ethosn_available(): + if not ethosn_available(): return def get_model(): @@ -84,7 +85,7 @@ def get_model(): def test_input_tuples(): - if not tei.ethosn_available(): + if not ethosn_available(): return def get_model(shapes, axis): From 01005f8cc1a51dfbea97dcc78cdee4680a8561a8 Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Mon, 10 Aug 2020 12:43:43 +0100 Subject: [PATCH 09/37] Remove unnecessary line Change-Id: Idf5cab853027adb0b0292de877e6dc02683821d7 --- src/relay/backend/contrib/ethosn/codegen_ethosn.h | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/src/relay/backend/contrib/ethosn/codegen_ethosn.h b/src/relay/backend/contrib/ethosn/codegen_ethosn.h index 72ed6ed86ecc..a955987278c4 100644 --- a/src/relay/backend/contrib/ethosn/codegen_ethosn.h +++ b/src/relay/backend/contrib/ethosn/codegen_ethosn.h @@ -292,8 +292,7 @@ class EthosnCompiler { cmms.emplace_back(CompileEthosnFunc(mod, name_node.value(), func)); } } else { - LOG(FATAL) << "The input ref is expected to be a Relay function or module" - << "\n"; + LOG(FATAL) << "The input ref is expected to be a Relay function"; } auto n = make_object(&cmms); return runtime::Module(n); From 175fcdd35a4cf858018e157e4dfe64c64be431a2 Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Mon, 10 Aug 2020 12:44:54 +0100 Subject: [PATCH 10/37] Simplify getting output_size Change-Id: If4643924768c2d7ea98525e9f792b7223cc2bcdf --- src/relay/backend/contrib/ethosn/codegen_ethosn.h | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/src/relay/backend/contrib/ethosn/codegen_ethosn.h b/src/relay/backend/contrib/ethosn/codegen_ethosn.h index a955987278c4..06a3fdd4b258 100644 --- a/src/relay/backend/contrib/ethosn/codegen_ethosn.h +++ b/src/relay/backend/contrib/ethosn/codegen_ethosn.h @@ -140,9 +140,8 @@ class InferTensorsVisitor : private ErrorReportingPass, private ExprVisitor { tensor_table_.clear(); CHECK(expr->checked_type().defined()); size_t output_size = 1; - if (expr->checked_type()->IsInstance()) { - auto type = expr->checked_type().as(); - output_size = type->fields.size(); + if (auto tuple = expr->checked_type().as()) { + output_size = tuple->fields.size(); } for (size_t i = 0; i < output_size; i++) { tensor_table_[expr].push_back(sl::TensorInfo({1, 1, 1, 1}, sl::DataType::UINT8_QUANTIZED, From 587e4a5c80c5df5fa446dc8a55ad08787f526749 Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Mon, 10 Aug 2020 12:46:38 +0100 Subject: [PATCH 11/37] Remove unnecessary new line Change-Id: Ia689c59cac28bd91e237ceecd829d8cf56d0d9c1 --- src/relay/backend/contrib/ethosn/ethosn_api.h | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/src/relay/backend/contrib/ethosn/ethosn_api.h b/src/relay/backend/contrib/ethosn/ethosn_api.h index ac3179504853..b61bdcf35b9e 100644 --- a/src/relay/backend/contrib/ethosn/ethosn_api.h +++ b/src/relay/backend/contrib/ethosn/ethosn_api.h @@ -121,8 +121,7 @@ class EthosnAPI { return EthosnError(); } - // Get a T from a constant represented by - // a NDArray. + // Get a T from a constant represented by a NDArray. template static EthosnError AsConstant(const Expr& expr, T* out) { if (!expr->IsInstance()) { From b8acaa6f5a7a94a9036f2f90654de190c956c0c9 Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Mon, 10 Aug 2020 12:51:21 +0100 Subject: [PATCH 12/37] Remove NOLINTS Change-Id: I149b97b28b516c7d9288a0858b2fbf1497e70250 --- src/runtime/contrib/ethosn/ethosn_runtime.h | 13 +++++++------ 1 file changed, 7 insertions(+), 6 deletions(-) diff --git a/src/runtime/contrib/ethosn/ethosn_runtime.h b/src/runtime/contrib/ethosn/ethosn_runtime.h index 6130518dbd0f..39cdba8cdacc 100644 --- a/src/runtime/contrib/ethosn/ethosn_runtime.h +++ b/src/runtime/contrib/ethosn/ethosn_runtime.h @@ -26,12 +26,13 @@ #include -#include -#include // NOLINT -#include // NOLINT -#include // NOLINT -#include // NOLINT -#include // NOLINT +#include +#include +#include +#include +#include + +#include "ethosn_support_library/Support.hpp" namespace tvm { namespace runtime { From 2f2a0544b443b27134cf377f6c50b8b7f6bf66ab Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Mon, 10 Aug 2020 13:15:12 +0100 Subject: [PATCH 13/37] Remove unused parts of PR Change-Id: I2db5b89d8fe2c114ab92305cdcf06d0fc45f4d2a --- .../backend/contrib/ethosn/ethosn_api.cc | 10 ----- src/relay/backend/contrib/ethosn/ethosn_api.h | 1 - .../contrib/test_ethosn/infrastructure.py | 39 ------------------- 3 files changed, 50 deletions(-) diff --git a/src/relay/backend/contrib/ethosn/ethosn_api.cc b/src/relay/backend/contrib/ethosn/ethosn_api.cc index 755ae6765fe8..ce453bf4a14f 100644 --- a/src/relay/backend/contrib/ethosn/ethosn_api.cc +++ b/src/relay/backend/contrib/ethosn/ethosn_api.cc @@ -130,16 +130,6 @@ sl::CompilationOptions EthosnAPI::CreateOptions() { return options; } -bool EthosnAPI::IsEthosFunc(const Call& call, const std::string& op_name) { - if (call->op->IsInstance()) { - Function func = Downcast(call->op); - CHECK(func.defined()); - auto name_node = func->GetAttr(attr::kComposite); - return name_node.value() == op_name; - } - return false; -} - bool EthosnAPI::IsEthosOp(const Call& call, const std::string& op_name) { if (call->op->IsInstance()) { Op op = Downcast(call->op); diff --git a/src/relay/backend/contrib/ethosn/ethosn_api.h b/src/relay/backend/contrib/ethosn/ethosn_api.h index b61bdcf35b9e..df3c62d8c879 100644 --- a/src/relay/backend/contrib/ethosn/ethosn_api.h +++ b/src/relay/backend/contrib/ethosn/ethosn_api.h @@ -92,7 +92,6 @@ class EthosnAPI { static sl::CompilationOptions CreateOptions(); - static bool IsEthosFunc(const Call& call, const std::string& op_name); static bool IsEthosOp(const Call& call, const std::string& op_name); static EthosnError Concatenate(const Expr& expr, ConcatenateParams* params); diff --git a/tests/python/contrib/test_ethosn/infrastructure.py b/tests/python/contrib/test_ethosn/infrastructure.py index 9dd33345c370..8f7e15c29a81 100644 --- a/tests/python/contrib/test_ethosn/infrastructure.py +++ b/tests/python/contrib/test_ethosn/infrastructure.py @@ -33,25 +33,6 @@ from tvm.relay.op.contrib import get_pattern_table -def get_real_image(im_height, im_width): - repo_base = 'https://github.com/dmlc/web-data/raw/master/tensorflow/models/InceptionV1/' - img_name = 'elephant-299.jpg' - image_url = os.path.join(repo_base, img_name) - img_path = download.download_testdata(image_url, img_name, module='data') - image = Image.open(img_path).resize((im_height, im_width)) - x = np.array(image).astype('uint8') - data = np.reshape(x, (1, im_height, im_width, 3)) - return data - - -def assert_lib_hash(lib, golden): - temp = util.tempdir() - path = temp.relpath("lib.cmm") - lib.imported_modules[1].save(path) - lib_hash = md5(open(path, 'rb').read()).hexdigest() - assert lib_hash == golden, "Expected hash: {} Got hash: {}".format(golden, lib_hash) - - def make_module(func, params): func = relay.Function(relay.analysis.free_vars(func), func) if len(params): @@ -174,26 +155,6 @@ def inference_result(checksum, outputs): return False -def generate_trials(space, r_factor=3): - np.random.seed(0) - max_len = 1 - for option in space: - max_len = max(max_len, len(option)) - - num_trials = r_factor * max_len - trials = [] - for i in range(num_trials): - trial = [] - for option in space: - if i % len(option) == 0: - np.random.shuffle(option) - trial.append(option[i % len(option)]) - - trials.append(trial) - - return trials - - def test_error(mod, params, err_msg): caught = None with tvm.transform.PassContext(opt_level=3): From 3eacad83142f38ad27b4428e073f7049a548bbcd Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Mon, 10 Aug 2020 13:15:53 +0100 Subject: [PATCH 14/37] Fix CI Ethos-N settings Change-Id: Idd955755d6f6d1cd3843462f627d0d952729e467 --- tests/scripts/task_config_build_cpu.sh | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tests/scripts/task_config_build_cpu.sh b/tests/scripts/task_config_build_cpu.sh index 89e6695b83f7..77b28e66fbb7 100755 --- a/tests/scripts/task_config_build_cpu.sh +++ b/tests/scripts/task_config_build_cpu.sh @@ -43,5 +43,5 @@ echo set\(USE_VTA_FSIM ON\) >> config.cmake echo set\(USE_TFLITE ON\) >> config.cmake echo set\(USE_TENSORFLOW_PATH \"/tensorflow\"\) >> config.cmake echo set\(USE_FLATBUFFERS_PATH \"/flatbuffers\"\) >> config.cmake -echo set\(USE_ETHOSN /opt/arm/ethosn-driver-dev/ethos-n77\) >> config.cmake -echo set\(USE_ETHOSN_HW ON\) >> config.cmake \ No newline at end of file +echo set\(USE_ETHOSN /opt/arm/ethosn-driver\) >> config.cmake +echo set\(USE_ETHOSN_HW OFF\) >> config.cmake \ No newline at end of file From dc48617ef906dae9a3d318e4928a533b4243933d Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Mon, 10 Aug 2020 13:17:31 +0100 Subject: [PATCH 15/37] Removed unnecessary line in infra Change-Id: I0ea866adf5d9166db85dd82d013a631d991ae633 --- tests/python/contrib/test_ethosn/infrastructure.py | 1 - 1 file changed, 1 deletion(-) diff --git a/tests/python/contrib/test_ethosn/infrastructure.py b/tests/python/contrib/test_ethosn/infrastructure.py index 8f7e15c29a81..fae8026e55c8 100644 --- a/tests/python/contrib/test_ethosn/infrastructure.py +++ b/tests/python/contrib/test_ethosn/infrastructure.py @@ -77,7 +77,6 @@ def __init__(self): def visit_call(self, call): if isinstance(call.op, tvm.ir.Op): self.count += 1 - super().visit_call(call) c = Counter() From 696a6d6d2f40a9f252a73ed5d939595fd190619f Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Mon, 10 Aug 2020 13:18:06 +0100 Subject: [PATCH 16/37] Remove unnecessary len in infra Change-Id: I869e8233d41c6ab7c2dc80f47d976c974043b80c --- tests/python/contrib/test_ethosn/infrastructure.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/python/contrib/test_ethosn/infrastructure.py b/tests/python/contrib/test_ethosn/infrastructure.py index fae8026e55c8..014fa594a2d3 100644 --- a/tests/python/contrib/test_ethosn/infrastructure.py +++ b/tests/python/contrib/test_ethosn/infrastructure.py @@ -35,7 +35,7 @@ def make_module(func, params): func = relay.Function(relay.analysis.free_vars(func), func) - if len(params): + if params: relay.build_module.bind_params_by_name(func, params) return tvm.IRModule.from_expr(func) From 80fb1083fb5f95903c907d8e19c611f70d3e1a24 Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Mon, 10 Aug 2020 13:21:45 +0100 Subject: [PATCH 17/37] Rename 'cpu_ops' to 'host_ops' Change-Id: I79a6ffcfd48cd055d279f493c672ec82f0c68e5c --- tests/python/contrib/test_ethosn/infrastructure.py | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/tests/python/contrib/test_ethosn/infrastructure.py b/tests/python/contrib/test_ethosn/infrastructure.py index 014fa594a2d3..257e16c124b4 100644 --- a/tests/python/contrib/test_ethosn/infrastructure.py +++ b/tests/python/contrib/test_ethosn/infrastructure.py @@ -68,7 +68,7 @@ def make_ethosn_partition(ethosn_expr): return mod -def get_cpu_op_count(mod): +def get_host_op_count(mod): class Counter(tvm.relay.ExprVisitor): def __init__(self): super().__init__() @@ -84,7 +84,7 @@ def visit_call(self, call): return c.count -def build(mod, params, npu=True, cpu_ops=0, npu_partitions=1): +def build(mod, params, npu=True, expected_host_ops=0, npu_partitions=1): relay.backend.compile_engine.get().clear() with tvm.transform.PassContext(opt_level=3, config={ "relay.ext.ethos-n.options": {"variant": 0} @@ -97,9 +97,9 @@ def build(mod, params, npu=True, cpu_ops=0, npu_partitions=1): mod = relay.transform.AnnotateTarget("ethos-n")(mod) mod = relay.transform.MergeCompilerRegions()(mod) mod = relay.transform.PartitionGraph()(mod) - cpu_op_count = get_cpu_op_count(mod) - assert cpu_op_count == cpu_ops, \ - "Got {} CPU operators, expected {}".format(cpu_op_count, cpu_ops) + host_op_count = get_host_op_count(mod) + assert host_op_count == expected_host_ops, \ + "Got {} host operators, expected {}".format(host_op_count, expected_host_ops) partition_count = 0 for global_var in mod.get_global_vars(): if "ethos-n" in global_var.name_hint: @@ -122,8 +122,8 @@ def run(graph, lib, params, inputs, outputs, npu=True): return out -def build_and_run(mod, inputs, outputs, params, ctx=tvm.cpu(), npu=True, cpu_ops=0, npu_partitions=1): - graph, lib, params = build(mod, params, npu, cpu_ops, npu_partitions) +def build_and_run(mod, inputs, outputs, params, ctx=tvm.cpu(), npu=True, expected_host_ops=0, npu_partitions=1): + graph, lib, params = build(mod, params, npu, expected_host_ops, npu_partitions) return run(graph, lib, params, inputs, outputs, npu) From df6e6d0cdc2071ec3a77178bc8aa4c7f890327eb Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Mon, 10 Aug 2020 13:25:40 +0100 Subject: [PATCH 18/37] Added explanation on mocking Change-Id: I1e88c07a47464e44cb45c6a327ec9c7e2d70cc94 --- src/runtime/contrib/ethosn/ethosn_device.cc | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/src/runtime/contrib/ethosn/ethosn_device.cc b/src/runtime/contrib/ethosn/ethosn_device.cc index 69139956c963..0ec6fa26a3fc 100644 --- a/src/runtime/contrib/ethosn/ethosn_device.cc +++ b/src/runtime/contrib/ethosn/ethosn_device.cc @@ -174,6 +174,11 @@ bool Inference(tvm::runtime::TVMArgs args, sl::CompiledNetwork* network, } // namespace tvm #else +/* If USE_ETHOSN_HW=OFF, we mock the inference call with a known-good output. + * That output can be set by using relay.ethos-n.test.infra.inference_result + * which will set the values the mocked inference will return the next time + * it's called. + */ #include #include From c81af106d3b857b73d042f0ec8a7d9f0a8aecf8f Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Mon, 10 Aug 2020 13:27:34 +0100 Subject: [PATCH 19/37] IsEthosOp -> IsEthosnOp Change-Id: I4fc1b462a74f8fae231ebafac614dd8d45be0feb --- src/relay/backend/contrib/ethosn/codegen.cc | 8 ++++---- src/relay/backend/contrib/ethosn/ethosn_api.cc | 2 +- src/relay/backend/contrib/ethosn/ethosn_api.h | 2 +- 3 files changed, 6 insertions(+), 6 deletions(-) diff --git a/src/relay/backend/contrib/ethosn/codegen.cc b/src/relay/backend/contrib/ethosn/codegen.cc index b477c0ce8cfc..a44eef5391dd 100644 --- a/src/relay/backend/contrib/ethosn/codegen.cc +++ b/src/relay/backend/contrib/ethosn/codegen.cc @@ -43,11 +43,11 @@ void InferTensorsVisitor::InferCall(const CallNode* cn) { EthosnError err; Call call = GetRef(cn); // Determine call -> NPU mapping - if (EthosnAPI::IsEthosOp(call, "qnn.concatenate")) { + if (EthosnAPI::IsEthosnOp(call, "qnn.concatenate")) { ConcatenateParams params; err = EthosnAPI::Concatenate(call, ¶ms); tensor_table_[cn->args[0]] = params.input_infos; - } else if (EthosnAPI::IsEthosOp(call, "split")) { + } else if (EthosnAPI::IsEthosnOp(call, "split")) { SplitParams params; params.input_info = GetTensorInfo(tensor_table_, call); err = EthosnAPI::Split(call, ¶ms); @@ -124,10 +124,10 @@ sl::TensorsAndId ConstructNetworkVisitor::HandleCall(const CallNode* cn) { sl::TensorAndId tensor; sl::TensorsAndId tensors; // Determine call -> NPU mapping - if (EthosnAPI::IsEthosOp(call, "qnn.concatenate")) { + if (EthosnAPI::IsEthosnOp(call, "qnn.concatenate")) { if ((err = MakeConcatenateLayer(call, &tensor))) ReportFatalError(call, err); return MakeOps(tensor); - } else if (EthosnAPI::IsEthosOp(call, "split")) { + } else if (EthosnAPI::IsEthosnOp(call, "split")) { if ((err = MakeSplitLayer(call, &tensors))) ReportFatalError(call, err); return tensors; } else { diff --git a/src/relay/backend/contrib/ethosn/ethosn_api.cc b/src/relay/backend/contrib/ethosn/ethosn_api.cc index ce453bf4a14f..cc879a38e591 100644 --- a/src/relay/backend/contrib/ethosn/ethosn_api.cc +++ b/src/relay/backend/contrib/ethosn/ethosn_api.cc @@ -130,7 +130,7 @@ sl::CompilationOptions EthosnAPI::CreateOptions() { return options; } -bool EthosnAPI::IsEthosOp(const Call& call, const std::string& op_name) { +bool EthosnAPI::IsEthosnOp(const Call& call, const std::string& op_name) { if (call->op->IsInstance()) { Op op = Downcast(call->op); CHECK(op.defined()); diff --git a/src/relay/backend/contrib/ethosn/ethosn_api.h b/src/relay/backend/contrib/ethosn/ethosn_api.h index df3c62d8c879..3e55d5116d26 100644 --- a/src/relay/backend/contrib/ethosn/ethosn_api.h +++ b/src/relay/backend/contrib/ethosn/ethosn_api.h @@ -92,7 +92,7 @@ class EthosnAPI { static sl::CompilationOptions CreateOptions(); - static bool IsEthosOp(const Call& call, const std::string& op_name); + static bool IsEthosnOp(const Call& call, const std::string& op_name); static EthosnError Concatenate(const Expr& expr, ConcatenateParams* params); static EthosnError Split(const Expr& expr, SplitParams* params); From d39524cfbcf6dd41c105dc43dc9659e96ba43cfa Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Mon, 10 Aug 2020 13:59:39 +0100 Subject: [PATCH 20/37] Improve documentation in ethosn_api.h Change-Id: I5586a7ba7ce71da667a6a9c6dd2e591028eb43b2 --- src/relay/backend/contrib/ethosn/codegen.cc | 18 ++++++-- .../backend/contrib/ethosn/ethosn_api.cc | 10 ----- src/relay/backend/contrib/ethosn/ethosn_api.h | 42 ++++++++++++++++++- 3 files changed, 54 insertions(+), 16 deletions(-) diff --git a/src/relay/backend/contrib/ethosn/codegen.cc b/src/relay/backend/contrib/ethosn/codegen.cc index a44eef5391dd..e0fa29c5a43a 100644 --- a/src/relay/backend/contrib/ethosn/codegen.cc +++ b/src/relay/backend/contrib/ethosn/codegen.cc @@ -39,15 +39,25 @@ sl::TensorInfo GetTensorInfo(std::map> tensor_ return sl::TensorInfo(); } +bool IsEthosnOp(const Call& call, const std::string& op_name) { + if (call->op->IsInstance()) { + Op op = Downcast(call->op); + CHECK(op.defined()); + return op == Op::Get(op_name); + } else { + return false; + } +} + void InferTensorsVisitor::InferCall(const CallNode* cn) { EthosnError err; Call call = GetRef(cn); // Determine call -> NPU mapping - if (EthosnAPI::IsEthosnOp(call, "qnn.concatenate")) { + if (IsEthosnOp(call, "qnn.concatenate")) { ConcatenateParams params; err = EthosnAPI::Concatenate(call, ¶ms); tensor_table_[cn->args[0]] = params.input_infos; - } else if (EthosnAPI::IsEthosnOp(call, "split")) { + } else if (IsEthosnOp(call, "split")) { SplitParams params; params.input_info = GetTensorInfo(tensor_table_, call); err = EthosnAPI::Split(call, ¶ms); @@ -124,10 +134,10 @@ sl::TensorsAndId ConstructNetworkVisitor::HandleCall(const CallNode* cn) { sl::TensorAndId tensor; sl::TensorsAndId tensors; // Determine call -> NPU mapping - if (EthosnAPI::IsEthosnOp(call, "qnn.concatenate")) { + if (IsEthosnOp(call, "qnn.concatenate")) { if ((err = MakeConcatenateLayer(call, &tensor))) ReportFatalError(call, err); return MakeOps(tensor); - } else if (EthosnAPI::IsEthosnOp(call, "split")) { + } else if (IsEthosnOp(call, "split")) { if ((err = MakeSplitLayer(call, &tensors))) ReportFatalError(call, err); return tensors; } else { diff --git a/src/relay/backend/contrib/ethosn/ethosn_api.cc b/src/relay/backend/contrib/ethosn/ethosn_api.cc index cc879a38e591..b469c0dc5f0a 100644 --- a/src/relay/backend/contrib/ethosn/ethosn_api.cc +++ b/src/relay/backend/contrib/ethosn/ethosn_api.cc @@ -130,16 +130,6 @@ sl::CompilationOptions EthosnAPI::CreateOptions() { return options; } -bool EthosnAPI::IsEthosnOp(const Call& call, const std::string& op_name) { - if (call->op->IsInstance()) { - Op op = Downcast(call->op); - CHECK(op.defined()); - return op == Op::Get(op_name); - } else { - return false; - } -} - EthosnError EthosnAPI::Concatenate(const Expr& expr, ConcatenateParams* params) { Call call = Downcast(expr); const auto& attrs = call->attrs.as(); diff --git a/src/relay/backend/contrib/ethosn/ethosn_api.h b/src/relay/backend/contrib/ethosn/ethosn_api.h index 3e55d5116d26..bd502a69acbe 100644 --- a/src/relay/backend/contrib/ethosn/ethosn_api.h +++ b/src/relay/backend/contrib/ethosn/ethosn_api.h @@ -55,6 +55,9 @@ struct SplitParams { sl::TensorInfo input_info; }; +/*! + * \brief A wrapper around std::stringstream to build an EthosnError. + */ class ErrStrm { public: template @@ -68,33 +71,68 @@ class ErrStrm { friend class EthosnError; }; +/*! + * \brief Custom error class for storing error messages produced + * during compilation for Ethos-N. + */ class EthosnError { public: + /*! \brief Default constructor */ EthosnError() {} + /*! + * \brief Construct error from an Array of Strings + * \param msgs The messages + */ explicit EthosnError(const Array& msgs) : msgs(msgs) {} + /*! + * \brief Construct error from a String + * \param msg The message + */ explicit EthosnError(const String& msg) { msgs.push_back(msg); } + /*! + * \brief Construct error from an ErrStrm + * \param err The ErrStrm + */ explicit EthosnError(const ErrStrm& err) : EthosnError(err.stream_.str()) {} + /*! \return Whether there are any error messages */ explicit operator bool() const { return !msgs.empty(); } + /*! \brief Add together two errors to give a single error with all the msgs */ EthosnError& operator+=(const EthosnError& other) { msgs.insert(msgs.end(), other.msgs.begin(), other.msgs.end()); return *this; } + /*! \brief The error messages */ Array msgs; }; +/*! + * \brief Functions to interact with Support Library's API including the + * translation of Relay ops/composite functions into Support Library + * equivalents. + */ class EthosnAPI { public: + /*! + * \brief Compile a Support Library network using the given compiler options + * \param network The network to be compiled + * \param options The options to compile with + * \return compiled_network The compiled network + */ static std::unique_ptr Compile(std::shared_ptr network, const sl::CompilationOptions& options); + /*! + * \brief Get the Support Library compilation options from the PassContext + * \return options The compilation options + */ static sl::CompilationOptions CreateOptions(); - static bool IsEthosnOp(const Call& call, const std::string& op_name); - + /*! \brief Extract the Support Library concatenate params from a Relay qnn.concatenate call */ static EthosnError Concatenate(const Expr& expr, ConcatenateParams* params); + /*! \brief Extract the Support Library split params from a Relay split call */ static EthosnError Split(const Expr& expr, SplitParams* params); private: From 3014543a8b280d083bd1556702150a32fc6ebec7 Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Mon, 10 Aug 2020 14:22:26 +0100 Subject: [PATCH 21/37] No longer iterate over module when compiling Change-Id: I80e1d494c6d574be06a2375e831343485712914d --- .../backend/contrib/ethosn/codegen_ethosn.h | 21 ++++++++----------- 1 file changed, 9 insertions(+), 12 deletions(-) diff --git a/src/relay/backend/contrib/ethosn/codegen_ethosn.h b/src/relay/backend/contrib/ethosn/codegen_ethosn.h index 06a3fdd4b258..f68bb75de706 100644 --- a/src/relay/backend/contrib/ethosn/codegen_ethosn.h +++ b/src/relay/backend/contrib/ethosn/codegen_ethosn.h @@ -259,18 +259,17 @@ NetworkWithIDs ConstructNetwork(const IRModule& mod, const GlobalVar& var, const class EthosnCompiler { public: static runtime::ethosn::OrderedCompiledNetwork CompileEthosnFunc(const IRModule& mod, - std::string name, + const GlobalVar& gvar, const Function& func) { // Construct the network - GlobalVar var = mod->GetGlobalVar(name); - auto network_with_ids = ConstructNetwork(mod, var, func); + auto network_with_ids = ConstructNetwork(mod, gvar, func); // Now set the required build flags sl::CompilationOptions options = EthosnAPI::CreateOptions(); // Finally compile the network auto compiled_network = EthosnAPI::Compile(network_with_ids.network, options); auto input_output_order = GetInputOutputOrder(network_with_ids, compiled_network); runtime::ethosn::OrderedCompiledNetwork ordered_network; - ordered_network.name = name; + ordered_network.name = gvar->name_hint; ordered_network.cmm = std::move(compiled_network); ordered_network.inputs = input_output_order.first; ordered_network.outputs = input_output_order.second; @@ -281,15 +280,13 @@ class EthosnCompiler { std::vector cmms; if (ref->IsInstance()) { IRModule mod; - Function bfunc = Downcast(ref); - auto name_node = bfunc->GetAttr(tvm::attr::kGlobalSymbol); + Function func = Downcast(ref); + auto name_node = func->GetAttr(tvm::attr::kGlobalSymbol); CHECK(name_node.defined()) << "Failed to retrieved external symbol."; - mod->Add(GlobalVar(name_node.value()), bfunc); - for (const auto& it : mod->functions) { - Function func = Downcast(it.second); - name_node = func->GetAttr(tvm::attr::kGlobalSymbol); - cmms.emplace_back(CompileEthosnFunc(mod, name_node.value(), func)); - } + GlobalVar gvar = GlobalVar(name_node.value()); + mod->Add(gvar, func); + Function mod_func = Downcast(mod->functions.at(gvar)); + cmms.emplace_back(CompileEthosnFunc(mod, gvar, mod_func)); } else { LOG(FATAL) << "The input ref is expected to be a Relay function"; } From c478165913d2d0f497e0c53c38546933749868ff Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Mon, 10 Aug 2020 14:28:14 +0100 Subject: [PATCH 22/37] Move EthosnCompiler implementations into codegen.cc Change-Id: I5bb6e9f62722d930d9dc040ac62bf87f29dd74c5 --- src/relay/backend/contrib/ethosn/codegen.cc | 53 +++++++++++++++++++ .../backend/contrib/ethosn/codegen_ethosn.h | 53 ++----------------- 2 files changed, 58 insertions(+), 48 deletions(-) diff --git a/src/relay/backend/contrib/ethosn/codegen.cc b/src/relay/backend/contrib/ethosn/codegen.cc index e0fa29c5a43a..93be45f8d278 100644 --- a/src/relay/backend/contrib/ethosn/codegen.cc +++ b/src/relay/backend/contrib/ethosn/codegen.cc @@ -218,6 +218,59 @@ EthosnError ConstructNetworkVisitor::MakeSplitLayer(const Call& call, sl::Tensor return EthosnError(); } +runtime::Module EthosnCompiler::CreateRuntimeModule(const ObjectRef& ref) { + std::vector cmms; + if (ref->IsInstance()) { + IRModule mod; + Function func = Downcast(ref); + auto name_node = func->GetAttr(tvm::attr::kGlobalSymbol); + CHECK(name_node.defined()) << "Failed to retrieved external symbol."; + GlobalVar gvar = GlobalVar(name_node.value()); + mod->Add(gvar, func); + Function mod_func = Downcast(mod->functions.at(gvar)); + cmms.emplace_back(CompileEthosnFunc(mod, gvar, mod_func)); + } else { + LOG(FATAL) << "The input ref is expected to be a Relay function"; + } + auto n = make_object(&cmms); + return runtime::Module(n); +} + +runtime::ethosn::OrderedCompiledNetwork EthosnCompiler::CompileEthosnFunc(const IRModule& mod, + const GlobalVar& gvar, + const Function& func) { + // Construct the network + auto network_with_ids = ConstructNetwork(mod, gvar, func); + // Now set the required build flags + sl::CompilationOptions options = EthosnAPI::CreateOptions(); + // Finally compile the network + auto compiled_network = EthosnAPI::Compile(network_with_ids.network, options); + auto input_output_order = GetInputOutputOrder(network_with_ids, compiled_network); + runtime::ethosn::OrderedCompiledNetwork ordered_network; + ordered_network.name = gvar->name_hint; + ordered_network.cmm = std::move(compiled_network); + ordered_network.inputs = input_output_order.first; + ordered_network.outputs = input_output_order.second; + return ordered_network; +} + +std::pair, std::vector> EthosnCompiler::GetInputOutputOrder( + NetworkWithIDs network, const std::unique_ptr& compiled_network) { + std::vector input_infos = compiled_network->GetInputBufferInfos(); + std::vector output_infos = compiled_network->GetOutputBufferInfos(); + std::vector input_order; + std::vector output_order; + for (const auto& input_info : input_infos) { + input_order.push_back(network.input_ids[input_info.m_SourceOperationId]); + } + for (const auto& output_info : output_infos) { + auto output_id = + std::make_pair(output_info.m_SourceOperationId, output_info.m_SourceOperationOutputIndex); + output_order.push_back(network.output_ids[output_id]); + } + return std::make_pair(input_order, output_order); +} + } // namespace ethosn } // namespace contrib } // namespace relay diff --git a/src/relay/backend/contrib/ethosn/codegen_ethosn.h b/src/relay/backend/contrib/ethosn/codegen_ethosn.h index f68bb75de706..e6322b6241cc 100644 --- a/src/relay/backend/contrib/ethosn/codegen_ethosn.h +++ b/src/relay/backend/contrib/ethosn/codegen_ethosn.h @@ -258,58 +258,15 @@ NetworkWithIDs ConstructNetwork(const IRModule& mod, const GlobalVar& var, const class EthosnCompiler { public: + static runtime::Module CreateRuntimeModule(const ObjectRef& ref); + + private: static runtime::ethosn::OrderedCompiledNetwork CompileEthosnFunc(const IRModule& mod, const GlobalVar& gvar, - const Function& func) { - // Construct the network - auto network_with_ids = ConstructNetwork(mod, gvar, func); - // Now set the required build flags - sl::CompilationOptions options = EthosnAPI::CreateOptions(); - // Finally compile the network - auto compiled_network = EthosnAPI::Compile(network_with_ids.network, options); - auto input_output_order = GetInputOutputOrder(network_with_ids, compiled_network); - runtime::ethosn::OrderedCompiledNetwork ordered_network; - ordered_network.name = gvar->name_hint; - ordered_network.cmm = std::move(compiled_network); - ordered_network.inputs = input_output_order.first; - ordered_network.outputs = input_output_order.second; - return ordered_network; - } - - static runtime::Module CreateRuntimeModule(const ObjectRef& ref) { - std::vector cmms; - if (ref->IsInstance()) { - IRModule mod; - Function func = Downcast(ref); - auto name_node = func->GetAttr(tvm::attr::kGlobalSymbol); - CHECK(name_node.defined()) << "Failed to retrieved external symbol."; - GlobalVar gvar = GlobalVar(name_node.value()); - mod->Add(gvar, func); - Function mod_func = Downcast(mod->functions.at(gvar)); - cmms.emplace_back(CompileEthosnFunc(mod, gvar, mod_func)); - } else { - LOG(FATAL) << "The input ref is expected to be a Relay function"; - } - auto n = make_object(&cmms); - return runtime::Module(n); - } + const Function& func); static std::pair, std::vector> GetInputOutputOrder( - NetworkWithIDs network, const std::unique_ptr& compiled_network) { - std::vector input_infos = compiled_network->GetInputBufferInfos(); - std::vector output_infos = compiled_network->GetOutputBufferInfos(); - std::vector input_order; - std::vector output_order; - for (const auto& input_info : input_infos) { - input_order.push_back(network.input_ids[input_info.m_SourceOperationId]); - } - for (const auto& output_info : output_infos) { - auto output_id = - std::make_pair(output_info.m_SourceOperationId, output_info.m_SourceOperationOutputIndex); - output_order.push_back(network.output_ids[output_id]); - } - return std::make_pair(input_order, output_order); - } + NetworkWithIDs network, const std::unique_ptr& compiled_network); }; runtime::Module CompileEthosn(const ObjectRef& ref) { From abb1066f01a547dae59e784ceaed553a5c70a7b3 Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Mon, 10 Aug 2020 14:28:58 +0100 Subject: [PATCH 23/37] Fix linting Change-Id: Ia44ec741a5330ad289cc6b5cd2bb1ed784fe6afc --- src/relay/backend/contrib/ethosn/codegen.cc | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/src/relay/backend/contrib/ethosn/codegen.cc b/src/relay/backend/contrib/ethosn/codegen.cc index 93be45f8d278..403e5ec4a15f 100644 --- a/src/relay/backend/contrib/ethosn/codegen.cc +++ b/src/relay/backend/contrib/ethosn/codegen.cc @@ -237,8 +237,8 @@ runtime::Module EthosnCompiler::CreateRuntimeModule(const ObjectRef& ref) { } runtime::ethosn::OrderedCompiledNetwork EthosnCompiler::CompileEthosnFunc(const IRModule& mod, - const GlobalVar& gvar, - const Function& func) { + const GlobalVar& gvar, + const Function& func) { // Construct the network auto network_with_ids = ConstructNetwork(mod, gvar, func); // Now set the required build flags @@ -255,7 +255,7 @@ runtime::ethosn::OrderedCompiledNetwork EthosnCompiler::CompileEthosnFunc(const } std::pair, std::vector> EthosnCompiler::GetInputOutputOrder( - NetworkWithIDs network, const std::unique_ptr& compiled_network) { + NetworkWithIDs network, const std::unique_ptr& compiled_network) { std::vector input_infos = compiled_network->GetInputBufferInfos(); std::vector output_infos = compiled_network->GetOutputBufferInfos(); std::vector input_order; @@ -265,7 +265,7 @@ std::pair, std::vector> EthosnCompiler::GetInput } for (const auto& output_info : output_infos) { auto output_id = - std::make_pair(output_info.m_SourceOperationId, output_info.m_SourceOperationOutputIndex); + std::make_pair(output_info.m_SourceOperationId, output_info.m_SourceOperationOutputIndex); output_order.push_back(network.output_ids[output_id]); } return std::make_pair(input_order, output_order); From 5cb6d9c2aadbcc1bd422ff2f0bd362863557b0f1 Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Mon, 10 Aug 2020 14:41:06 +0100 Subject: [PATCH 24/37] Refactor EthosnAPI compilation functions into EthosnCompiler Change-Id: Iee0aecbe43a84fefb437ab9ff064e3f8b42c80a4 --- src/relay/backend/contrib/ethosn/codegen.cc | 42 ++++++++- .../backend/contrib/ethosn/codegen_ethosn.h | 87 ++++++++++++++++++ .../backend/contrib/ethosn/ethosn_api.cc | 90 ------------------- src/relay/backend/contrib/ethosn/ethosn_api.h | 15 ---- 4 files changed, 127 insertions(+), 107 deletions(-) diff --git a/src/relay/backend/contrib/ethosn/codegen.cc b/src/relay/backend/contrib/ethosn/codegen.cc index 403e5ec4a15f..ec41440b37c2 100644 --- a/src/relay/backend/contrib/ethosn/codegen.cc +++ b/src/relay/backend/contrib/ethosn/codegen.cc @@ -24,6 +24,7 @@ #include #include +#include "capabilities.h" #include "codegen_ethosn.h" #include "ethosn_api.h" @@ -242,10 +243,17 @@ runtime::ethosn::OrderedCompiledNetwork EthosnCompiler::CompileEthosnFunc(const // Construct the network auto network_with_ids = ConstructNetwork(mod, gvar, func); // Now set the required build flags - sl::CompilationOptions options = EthosnAPI::CreateOptions(); + sl::CompilationOptions options = CreateOptions(); // Finally compile the network - auto compiled_network = EthosnAPI::Compile(network_with_ids.network, options); + std::vector> compiled_networks = + sl::Compile(*network_with_ids.network, options); + CHECK_GE(compiled_networks.size(), 1) << "Ethos-N compiler failed to compile network"; + auto compiled_network = std::move(compiled_networks[0]); + // Determine the order that the inputs/outputs are in and how that corresponds to the + // order that the TVM runtime will expect them in auto input_output_order = GetInputOutputOrder(network_with_ids, compiled_network); + // Use the order information to create an 'ordered' network with includes how to map + // the inputs/outputs from the TVM runtime to the inputs/outputs of the compiled network runtime::ethosn::OrderedCompiledNetwork ordered_network; ordered_network.name = gvar->name_hint; ordered_network.cmm = std::move(compiled_network); @@ -254,15 +262,45 @@ runtime::ethosn::OrderedCompiledNetwork EthosnCompiler::CompileEthosnFunc(const return ordered_network; } +sl::CompilationOptions EthosnCompiler::CreateOptions() { + auto ctx = transform::PassContext::Current(); + auto cfg = ctx->GetConfig("relay.ext.ethos-n.options"); + if (!cfg.defined()) { + cfg = AttrsWithDefaultValues(); + } + + sl::CompilationOptions options(variants[cfg.value()->variant]); + options.m_Strategy0 = cfg.value()->strategy0; + options.m_Strategy1 = cfg.value()->strategy1; + options.m_Strategy3 = cfg.value()->strategy3; + options.m_Strategy4 = cfg.value()->strategy4; + options.m_Strategy6 = cfg.value()->strategy6; + options.m_Strategy7 = cfg.value()->strategy7; + options.m_DebugInfo.m_DumpRam = cfg.value()->dump_ram; + options.m_DebugInfo.m_InitialSramDump = cfg.value()->initial_sram_dump; + options.m_BlockConfig16x16 = cfg.value()->block_config_16x16; + options.m_BlockConfig32x8 = cfg.value()->block_config_32x8; + options.m_BlockConfig8x32 = cfg.value()->block_config_8x32; + options.m_BlockConfig8x8 = cfg.value()->block_config_8x8; + options.m_EnableIntermediateCompression = cfg.value()->enable_intermediate_compression; + options.m_DisableWinograd = cfg.value()->disable_winograd; + options.m_DebugInfo.m_DumpDebugFiles = cfg.value()->dump_debug_files; + options.m_DebugInfo.m_DebugDir = cfg.value()->debug_dir; + options.m_EnableCascading = cfg.value()->enable_cascading; + return options; +} + std::pair, std::vector> EthosnCompiler::GetInputOutputOrder( NetworkWithIDs network, const std::unique_ptr& compiled_network) { std::vector input_infos = compiled_network->GetInputBufferInfos(); std::vector output_infos = compiled_network->GetOutputBufferInfos(); std::vector input_order; std::vector output_order; + // Find the order of the inputs in the compiled network for (const auto& input_info : input_infos) { input_order.push_back(network.input_ids[input_info.m_SourceOperationId]); } + // Find the order of the outputs in the compiled network for (const auto& output_info : output_infos) { auto output_id = std::make_pair(output_info.m_SourceOperationId, output_info.m_SourceOperationOutputIndex); diff --git a/src/relay/backend/contrib/ethosn/codegen_ethosn.h b/src/relay/backend/contrib/ethosn/codegen_ethosn.h index e6322b6241cc..6e248aa43513 100644 --- a/src/relay/backend/contrib/ethosn/codegen_ethosn.h +++ b/src/relay/backend/contrib/ethosn/codegen_ethosn.h @@ -256,15 +256,102 @@ NetworkWithIDs ConstructNetwork(const IRModule& mod, const GlobalVar& var, const return ConstructNetworkVisitor(mod, var).Construct(func); } +/*! \brief Attributes to store the compiler options for Ethos-N */ +struct EthosnCompilerConfigNode : public tvm::AttrsNode { + int variant; + bool strategy0; + bool strategy1; + bool strategy3; + bool strategy4; + bool strategy6; + bool strategy7; + bool dump_ram; + bool initial_sram_dump; + bool block_config_16x16; + bool block_config_32x8; + bool block_config_8x32; + bool block_config_8x8; + bool enable_intermediate_compression; + bool disable_winograd; + bool dump_debug_files; + String debug_dir; + bool enable_cascading; + + TVM_DECLARE_ATTRS(EthosnCompilerConfigNode, "ext.attrs.EthosnCompilerConfigNode") { + TVM_ATTR_FIELD(variant) + .describe("0 for Ethos-N77, 1 for Ethos-N57, 2 for Ethos-N37. See Ethos-N documentation.") + .set_default(0); + TVM_ATTR_FIELD(strategy0).set_default(true); + TVM_ATTR_FIELD(strategy1).set_default(true); + TVM_ATTR_FIELD(strategy3).set_default(true); + TVM_ATTR_FIELD(strategy4).set_default(true); + TVM_ATTR_FIELD(strategy6).set_default(true); + TVM_ATTR_FIELD(strategy7).set_default(true); + TVM_ATTR_FIELD(dump_ram).set_default(false); + TVM_ATTR_FIELD(initial_sram_dump).set_default(false); + TVM_ATTR_FIELD(block_config_16x16).set_default(true); + TVM_ATTR_FIELD(block_config_32x8).set_default(true); + TVM_ATTR_FIELD(block_config_8x32).set_default(true); + TVM_ATTR_FIELD(block_config_8x8).set_default(true); + TVM_ATTR_FIELD(enable_intermediate_compression).set_default(true); + TVM_ATTR_FIELD(disable_winograd).set_default(false); + TVM_ATTR_FIELD(dump_debug_files).set_default(false); + TVM_ATTR_FIELD(debug_dir).set_default("."); + TVM_ATTR_FIELD(enable_cascading).set_default(false); + } +}; + +class EthosnCompilerConfig : public Attrs { + public: + TVM_DEFINE_NOTNULLABLE_OBJECT_REF_METHODS(EthosnCompilerConfig, Attrs, EthosnCompilerConfigNode); +}; + +TVM_REGISTER_NODE_TYPE(EthosnCompilerConfigNode); +TVM_REGISTER_PASS_CONFIG_OPTION("relay.ext.ethos-n.options", EthosnCompilerConfig); + +/*! \brief The compiler for Ethos-N functions */ class EthosnCompiler { public: + /*! + * \brief Create an Ethos-N runtime module from a Relay Ethos-N function + * \param ref An ObjectRef pointing to a Relay Ethos-N function + * \return runtime_module An Ethos-N runtime module + */ static runtime::Module CreateRuntimeModule(const ObjectRef& ref); private: + /*! + * \brief Compile a single Relay Ethos-N function into an ordered compiled network. + * Compilation options will be taken from the PassContext. + * \param mod The module the function is stored in (for error reporting purposes) + * \param gvar The global var corresponding to the function + * \param func The function to be compiled + * \return ordered_compiled_network A compiled network with additional information + * to handle difference in input/output ordering between the TVM runtime and the + * Ethos-N compiled network. + */ static runtime::ethosn::OrderedCompiledNetwork CompileEthosnFunc(const IRModule& mod, const GlobalVar& gvar, const Function& func); + /*! + * \brief Get the Support Library compilation options from the PassContext + * \return options The compilation options + */ + static sl::CompilationOptions CreateOptions(); + + /*! + * \brief Determine the order in which inputs should be provided/outputs should be + * read from a compiled network. This is required because when you compile a network + * for Ethos-N, you don't have control over the order in which the inputs/outputs + * are given. You can, however, query what order the compiler decided to give them in. + * We therefore keep track of our desired order and the actual order and create a + * small translation table between the two for use in the runtime. + * \param network A network additionally with the desired input/output order + * \param compiled_network The compiled network with an as yet undetermined input/output order + * \return input_output_order The order in which to permute the inputs/outputs given + * by the TVM runtime such that they map correctly to the compiled network. + */ static std::pair, std::vector> GetInputOutputOrder( NetworkWithIDs network, const std::unique_ptr& compiled_network); }; diff --git a/src/relay/backend/contrib/ethosn/ethosn_api.cc b/src/relay/backend/contrib/ethosn/ethosn_api.cc index b469c0dc5f0a..6e85cd7e3f4d 100644 --- a/src/relay/backend/contrib/ethosn/ethosn_api.cc +++ b/src/relay/backend/contrib/ethosn/ethosn_api.cc @@ -32,7 +32,6 @@ #include #include -#include "capabilities.h" #include "ethosn_support_library/Support.hpp" #include "ethosn_support_library/SupportQueries.hpp" @@ -41,95 +40,6 @@ namespace relay { namespace contrib { namespace ethosn { -std::unique_ptr EthosnAPI::Compile(std::shared_ptr network, - const sl::CompilationOptions& options) { - std::vector> compiled_network = - sl::Compile(*network, options); - CHECK_GE(compiled_network.size(), 1) << "Ethos-N compiler failed to compile network"; - - return std::move(compiled_network[0]); -} - -struct EthosnCompilerConfigNode : public tvm::AttrsNode { - int variant; - bool strategy0; - bool strategy1; - bool strategy3; - bool strategy4; - bool strategy6; - bool strategy7; - bool dump_ram; - bool initial_sram_dump; - bool block_config_16x16; - bool block_config_32x8; - bool block_config_8x32; - bool block_config_8x8; - bool enable_intermediate_compression; - bool disable_winograd; - bool dump_debug_files; - String debug_dir; - bool enable_cascading; - - TVM_DECLARE_ATTRS(EthosnCompilerConfigNode, "ext.attrs.EthosnCompilerConfigNode") { - TVM_ATTR_FIELD(variant) - .describe("0 for Ethos-N77, 1 for Ethos-N57, 2 for Ethos-N37. See Ethos-N documentation.") - .set_default(0); - TVM_ATTR_FIELD(strategy0).set_default(true); - TVM_ATTR_FIELD(strategy1).set_default(true); - TVM_ATTR_FIELD(strategy3).set_default(true); - TVM_ATTR_FIELD(strategy4).set_default(true); - TVM_ATTR_FIELD(strategy6).set_default(true); - TVM_ATTR_FIELD(strategy7).set_default(true); - TVM_ATTR_FIELD(dump_ram).set_default(false); - TVM_ATTR_FIELD(initial_sram_dump).set_default(false); - TVM_ATTR_FIELD(block_config_16x16).set_default(true); - TVM_ATTR_FIELD(block_config_32x8).set_default(true); - TVM_ATTR_FIELD(block_config_8x32).set_default(true); - TVM_ATTR_FIELD(block_config_8x8).set_default(true); - TVM_ATTR_FIELD(enable_intermediate_compression).set_default(true); - TVM_ATTR_FIELD(disable_winograd).set_default(false); - TVM_ATTR_FIELD(dump_debug_files).set_default(false); - TVM_ATTR_FIELD(debug_dir).set_default("."); - TVM_ATTR_FIELD(enable_cascading).set_default(false); - } -}; - -class EthosnCompilerConfig : public Attrs { - public: - TVM_DEFINE_NOTNULLABLE_OBJECT_REF_METHODS(EthosnCompilerConfig, Attrs, EthosnCompilerConfigNode); -}; - -TVM_REGISTER_NODE_TYPE(EthosnCompilerConfigNode); -TVM_REGISTER_PASS_CONFIG_OPTION("relay.ext.ethos-n.options", EthosnCompilerConfig); - -sl::CompilationOptions EthosnAPI::CreateOptions() { - auto ctx = transform::PassContext::Current(); - auto cfg = ctx->GetConfig("relay.ext.ethos-n.options"); - if (!cfg.defined()) { - cfg = AttrsWithDefaultValues(); - } - - sl::CompilationOptions options(variants[cfg.value()->variant]); - options.m_Strategy0 = cfg.value()->strategy0; - options.m_Strategy1 = cfg.value()->strategy1; - options.m_Strategy3 = cfg.value()->strategy3; - options.m_Strategy4 = cfg.value()->strategy4; - options.m_Strategy6 = cfg.value()->strategy6; - options.m_Strategy7 = cfg.value()->strategy7; - options.m_DebugInfo.m_DumpRam = cfg.value()->dump_ram; - options.m_DebugInfo.m_InitialSramDump = cfg.value()->initial_sram_dump; - options.m_BlockConfig16x16 = cfg.value()->block_config_16x16; - options.m_BlockConfig32x8 = cfg.value()->block_config_32x8; - options.m_BlockConfig8x32 = cfg.value()->block_config_8x32; - options.m_BlockConfig8x8 = cfg.value()->block_config_8x8; - options.m_EnableIntermediateCompression = cfg.value()->enable_intermediate_compression; - options.m_DisableWinograd = cfg.value()->disable_winograd; - options.m_DebugInfo.m_DumpDebugFiles = cfg.value()->dump_debug_files; - options.m_DebugInfo.m_DebugDir = cfg.value()->debug_dir; - options.m_EnableCascading = cfg.value()->enable_cascading; - return options; -} - EthosnError EthosnAPI::Concatenate(const Expr& expr, ConcatenateParams* params) { Call call = Downcast(expr); const auto& attrs = call->attrs.as(); diff --git a/src/relay/backend/contrib/ethosn/ethosn_api.h b/src/relay/backend/contrib/ethosn/ethosn_api.h index bd502a69acbe..0f9bc01ba6be 100644 --- a/src/relay/backend/contrib/ethosn/ethosn_api.h +++ b/src/relay/backend/contrib/ethosn/ethosn_api.h @@ -115,21 +115,6 @@ class EthosnError { */ class EthosnAPI { public: - /*! - * \brief Compile a Support Library network using the given compiler options - * \param network The network to be compiled - * \param options The options to compile with - * \return compiled_network The compiled network - */ - static std::unique_ptr Compile(std::shared_ptr network, - const sl::CompilationOptions& options); - - /*! - * \brief Get the Support Library compilation options from the PassContext - * \return options The compilation options - */ - static sl::CompilationOptions CreateOptions(); - /*! \brief Extract the Support Library concatenate params from a Relay qnn.concatenate call */ static EthosnError Concatenate(const Expr& expr, ConcatenateParams* params); /*! \brief Extract the Support Library split params from a Relay split call */ From 186c92e6e4301f87d22e99929917255874a0689e Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Mon, 10 Aug 2020 14:58:30 +0100 Subject: [PATCH 25/37] Improve docs for Tvm2Npu Change-Id: Ia39e9e1508513ca39c1d585fbccc3ae38fcbb9fb --- src/relay/backend/contrib/ethosn/ethosn_api.h | 2 ++ 1 file changed, 2 insertions(+) diff --git a/src/relay/backend/contrib/ethosn/ethosn_api.h b/src/relay/backend/contrib/ethosn/ethosn_api.h index 0f9bc01ba6be..d1bc120ef7f7 100644 --- a/src/relay/backend/contrib/ethosn/ethosn_api.h +++ b/src/relay/backend/contrib/ethosn/ethosn_api.h @@ -121,7 +121,9 @@ class EthosnAPI { static EthosnError Split(const Expr& expr, SplitParams* params); private: + /*! \brief Convert a TVM tensor shape to a SL tensor shape */ static EthosnError Tvm2Npu(const Array& shape, sl::TensorShape* npu_shape); + /*! \brief Convert a TVM data type to a SL data type */ static EthosnError Tvm2Npu(const tvm::DataType& dtype, sl::DataType* data_type); // Convert an array of IntImmNodes into ValueT From 6a967e78cbff6b2809db9e7b9cca85c180e4338a Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Mon, 10 Aug 2020 15:03:05 +0100 Subject: [PATCH 26/37] Move more implementation out of headers Change-Id: I1e33084ceb520b75f06b4d7a4acff5b9b2225bd5 --- src/relay/backend/contrib/ethosn/codegen.cc | 46 +++++++++++++++++++ .../backend/contrib/ethosn/codegen_ethosn.h | 46 +------------------ 2 files changed, 48 insertions(+), 44 deletions(-) diff --git a/src/relay/backend/contrib/ethosn/codegen.cc b/src/relay/backend/contrib/ethosn/codegen.cc index ec41440b37c2..f66eb94cb20f 100644 --- a/src/relay/backend/contrib/ethosn/codegen.cc +++ b/src/relay/backend/contrib/ethosn/codegen.cc @@ -50,6 +50,21 @@ bool IsEthosnOp(const Call& call, const std::string& op_name) { } } +std::map> InferTensorsVisitor::Infer(const Expr& expr) { + tensor_table_.clear(); + CHECK(expr->checked_type().defined()); + size_t output_size = 1; + if (auto tuple = expr->checked_type().as()) { + output_size = tuple->fields.size(); + } + for (size_t i = 0; i < output_size; i++) { + tensor_table_[expr].push_back(sl::TensorInfo({1, 1, 1, 1}, sl::DataType::UINT8_QUANTIZED, + sl::DataFormat::NHWC, sl::QuantizationInfo())); + } + VisitInferred(expr); + return tensor_table_; +} + void InferTensorsVisitor::InferCall(const CallNode* cn) { EthosnError err; Call call = GetRef(cn); @@ -129,6 +144,37 @@ sl::TensorsAndId MakeOps(const sl::TensorAndId& op) { return ops; } +NetworkWithIDs ConstructNetworkVisitor::Construct(const Function& func) { + // Initialise everything + NetworkWithIDs network_with_ids; + network_ = sl::CreateNetwork(); + network_with_ids.network = network_; + operand_table_.clear(); + + // Infer tensor information + tensor_table_ = InferTensors(this->mod_, this->var_, func->body); + // Add the inputs in the order they appear in the parameters + unsigned int idx = 0; + for (const auto& param : func->params) { + for (const auto& tensor_info : tensor_table_[param]) { + auto tensor_and_id = AddInput(network_, tensor_info); + operand_table_[param].push_back(tensor_and_id.tensor); + id_table_[param].push_back(std::make_pair(tensor_and_id.operationId, 0)); + network_with_ids.input_ids[tensor_and_id.operationId] = idx++; + } + } + // Add the function body + VisitExpr(func->body); + // Add the outputs + idx = 0; + for (const auto& layer : operand_table_[func->body]) { + AddOutput(network_, *layer); + network_with_ids.output_ids[id_table_[func->body][idx]] = idx; + idx++; + } + return network_with_ids; +} + sl::TensorsAndId ConstructNetworkVisitor::HandleCall(const CallNode* cn) { EthosnError err; Call call = GetRef(cn); diff --git a/src/relay/backend/contrib/ethosn/codegen_ethosn.h b/src/relay/backend/contrib/ethosn/codegen_ethosn.h index 6e248aa43513..714a22d22027 100644 --- a/src/relay/backend/contrib/ethosn/codegen_ethosn.h +++ b/src/relay/backend/contrib/ethosn/codegen_ethosn.h @@ -136,20 +136,7 @@ class InferTensorsVisitor : private ErrorReportingPass, private ExprVisitor { * \note This algorithm does not traverse into functions, so call it on * the body of the function you're interested in. */ - std::map> Infer(const Expr& expr) { - tensor_table_.clear(); - CHECK(expr->checked_type().defined()); - size_t output_size = 1; - if (auto tuple = expr->checked_type().as()) { - output_size = tuple->fields.size(); - } - for (size_t i = 0; i < output_size; i++) { - tensor_table_[expr].push_back(sl::TensorInfo({1, 1, 1, 1}, sl::DataType::UINT8_QUANTIZED, - sl::DataFormat::NHWC, sl::QuantizationInfo())); - } - VisitInferred(expr); - return tensor_table_; - } + std::map> Infer(const Expr& expr); private: // Infer a callnode if it's a supported operator/composite function @@ -198,36 +185,7 @@ class ConstructNetworkVisitor : public MixedModeVisitor, private ErrorReportingP * \return A support library network that performs the same operation as the Relay * function. */ - NetworkWithIDs Construct(const Function& func) { - // Initialise everything - NetworkWithIDs network_with_ids; - network_ = sl::CreateNetwork(); - network_with_ids.network = network_; - operand_table_.clear(); - - // Infer tensor information - tensor_table_ = InferTensors(this->mod_, this->var_, func->body); - // Add the inputs in the order they appear in the parameters - unsigned int idx = 0; - for (const auto& param : func->params) { - for (const auto& tensor_info : tensor_table_[param]) { - auto tensor_and_id = AddInput(network_, tensor_info); - operand_table_[param].push_back(tensor_and_id.tensor); - id_table_[param].push_back(std::make_pair(tensor_and_id.operationId, 0)); - network_with_ids.input_ids[tensor_and_id.operationId] = idx++; - } - } - // Add the function body - VisitExpr(func->body); - // Add the outputs - idx = 0; - for (const auto& layer : operand_table_[func->body]) { - AddOutput(network_, *layer); - network_with_ids.output_ids[id_table_[func->body][idx]] = idx; - idx++; - } - return network_with_ids; - } + NetworkWithIDs Construct(const Function& func); private: // Translate from a callnode to the appropriate 'Make' method From c1b5b96074aedf13807da38323dd5c6a997d1322 Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Mon, 10 Aug 2020 15:25:39 +0100 Subject: [PATCH 27/37] Move implementation in ethosn_api.h Change-Id: I51ab386892a2aa84aa47d03641aac8468f5737ae --- .../backend/contrib/ethosn/ethosn_api.cc | 30 +++++++++++++++++++ src/relay/backend/contrib/ethosn/ethosn_api.h | 24 ++------------- 2 files changed, 32 insertions(+), 22 deletions(-) diff --git a/src/relay/backend/contrib/ethosn/ethosn_api.cc b/src/relay/backend/contrib/ethosn/ethosn_api.cc index 6e85cd7e3f4d..cbd10bb9d743 100644 --- a/src/relay/backend/contrib/ethosn/ethosn_api.cc +++ b/src/relay/backend/contrib/ethosn/ethosn_api.cc @@ -128,6 +128,36 @@ EthosnError EthosnAPI::Tvm2Npu(const tvm::DataType& dtype, sl::DataType* data_ty return EthosnError(ErrStrm() << "dtype=\'" << dtype << "\', dtype must be either uint8 or int32"); } +// Convert an array of IntImmNodes into ValueT +// IndexT type of Array indexing variable +// ValueT type of resulting value +template +EthosnError EthosnAPI::AsArray(const Array& arr, std::array* v) { + if (arr.size() > 4) + return EthosnError(ErrStrm() << "dimensions=" << arr.size() << ", dimensions must be <= 4"); + for (size_t i = 0; i < std::min(arr.size(), 4ul); i++) { + const PrimExpr& a = arr[i]; + const auto* intImm = a.as(); + if (intImm->value > std::numeric_limits::max()) { + return EthosnError(ErrStrm() << "axis size=" << intImm->value << ", axis size must be <= " + << std::numeric_limits::max()); + } + (*v)[i] = static_cast(intImm->value); + } + return EthosnError(); +} + +// Get a T from a constant represented by a NDArray. +template +EthosnError EthosnAPI::AsConstant(const Expr& expr, T* out) { + if (!expr->IsInstance()) { + return EthosnError("expected constant data"); + } + runtime::NDArray data = Downcast(expr)->data; + *out = *static_cast(data.operator->()->data); + return EthosnError(); +} + TVM_REGISTER_GLOBAL("relay.ethos-n.support.concatenate") .set_body([](tvm::TVMArgs args, tvm::TVMRetValue* rv) { Call call = args[0]; diff --git a/src/relay/backend/contrib/ethosn/ethosn_api.h b/src/relay/backend/contrib/ethosn/ethosn_api.h index d1bc120ef7f7..34af7ce0b1d8 100644 --- a/src/relay/backend/contrib/ethosn/ethosn_api.h +++ b/src/relay/backend/contrib/ethosn/ethosn_api.h @@ -130,31 +130,11 @@ class EthosnAPI { // IndexT type of Array indexing variable // ValueT type of resulting value template - static EthosnError AsArray(const Array& arr, std::array* v) { - if (arr.size() > 4) - return EthosnError(ErrStrm() << "dimensions=" << arr.size() << ", dimensions must be <= 4"); - for (size_t i = 0; i < std::min(arr.size(), 4ul); i++) { - const PrimExpr& a = arr[i]; - const auto* intImm = a.as(); - if (intImm->value > std::numeric_limits::max()) { - return EthosnError(ErrStrm() << "axis size=" << intImm->value << ", axis size must be <= " - << std::numeric_limits::max()); - } - (*v)[i] = static_cast(intImm->value); - } - return EthosnError(); - } + static EthosnError AsArray(const Array& arr, std::array* v); // Get a T from a constant represented by a NDArray. template - static EthosnError AsConstant(const Expr& expr, T* out) { - if (!expr->IsInstance()) { - return EthosnError("expected constant data"); - } - runtime::NDArray data = Downcast(expr)->data; - *out = *static_cast(data.operator->()->data); - return EthosnError(); - } + static EthosnError AsConstant(const Expr& expr, T* out); }; } // namespace ethosn From d12983d3c4db8746970215e34d9cb69650620e22 Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Mon, 10 Aug 2020 12:56:41 +0100 Subject: [PATCH 28/37] Improve docs for capabilities.h Change-Id: Iaaee508aafa1cbb7650a04ed87bd6c1b91823a58 --- src/relay/backend/contrib/ethosn/capabilities.h | 14 ++++++++++++++ 1 file changed, 14 insertions(+) diff --git a/src/relay/backend/contrib/ethosn/capabilities.h b/src/relay/backend/contrib/ethosn/capabilities.h index 71eb13be4691..409d440cf8e1 100644 --- a/src/relay/backend/contrib/ethosn/capabilities.h +++ b/src/relay/backend/contrib/ethosn/capabilities.h @@ -17,6 +17,15 @@ * under the License. */ +/*! + * \file src/relay/backend/contrib/ethosn/capabilities.h + * \brief The Ethos-N processor series has four variants, the Ethos-N37, Ethos-N57, Ethos-N77 + * and the Ethos-N78. This release of the integration supports the first three variants. + * Configuration information for each variant is stored as a blob in this file. These blobs + * are passed into the Ethos-N support library, which in turn uses them to optimize the + * generated command-stream appropriately for the specified variant. + */ + #ifndef TVM_RELAY_BACKEND_CONTRIB_ETHOSN_CAPABILITIES_H_ #define TVM_RELAY_BACKEND_CONTRIB_ETHOSN_CAPABILITIES_H_ @@ -27,6 +36,11 @@ namespace relay { namespace contrib { namespace ethosn { +/* Ethos-N variants (N77, N57 and N37) + * variant[0] - N77 + * variant[1] - N57 + * variant[2] - N37 + */ static std::vector variants[3] = { { 0x02, 0x00, 0x00, 0x00, 0x74, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, From 93c5091ce5e1224625050b44dd61fc41eca1e8ca Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Wed, 12 Aug 2020 15:20:42 +0100 Subject: [PATCH 29/37] Use else() in cmake Change-Id: I4b64a87f32b3616ec87c9937d9fc998b8dc5d7b4 --- cmake/modules/contrib/EthosN.cmake | 5 ++--- cmake/util/FindEthosN.cmake | 5 ++--- 2 files changed, 4 insertions(+), 6 deletions(-) diff --git a/cmake/modules/contrib/EthosN.cmake b/cmake/modules/contrib/EthosN.cmake index 7c694822b21e..ca1f7daa8c8a 100644 --- a/cmake/modules/contrib/EthosN.cmake +++ b/cmake/modules/contrib/EthosN.cmake @@ -22,9 +22,8 @@ if(NOT USE_ETHOSN STREQUAL "OFF") if(NOT ETHOSN_FOUND) message(FATAL_ERROR "Cannot find Ethos-N, USE_ETHOSN=" ${USE_ETHOSN}) - endif() - if (ETHOSN_FOUND) + else() include_directories(${ETHOSN_INCLUDE_DIRS}) add_definitions(${ETHOSN_DEFINITIONS}) @@ -50,7 +49,7 @@ if(NOT USE_ETHOSN STREQUAL "OFF") set_source_files_properties(${COMPILER_ETHOSN_SRCS} PROPERTIES COMPILE_FLAGS "-fno-rtti") endif() - endif(ETHOSN_FOUND) + endif(NOT ETHOSN_FOUND) else() if(USE_ETHOSN_HW) message(FATAL_ERROR "Cannot enable Ethos-N HW if USE_ETHOSN=OFF") diff --git a/cmake/util/FindEthosN.cmake b/cmake/util/FindEthosN.cmake index 7e8c7b90b756..7f0fb6499691 100644 --- a/cmake/util/FindEthosN.cmake +++ b/cmake/util/FindEthosN.cmake @@ -82,14 +82,13 @@ macro(find_ethosn use_ethosn) if(__use_ethosn STREQUAL "ON") message(WARNING "No cmake find_package available for Arm Ethos-N") endif() - endif() # additional libraries - if(ETHOSN_FOUND) + else() message(STATUS "Found ETHOSN_DEFINITIONS=${ETHOSN_DEFINITIONS}") message(STATUS "Found ETHOSN_INCLUDE_DIRS=${ETHOSN_INCLUDE_DIRS}") message(STATUS "Found ETHOSN_COMPILER_LIBRARY=${ETHOSN_COMPILER_LIBRARY}") message(STATUS "Found ETHOSN_RUNTIME_LIBRARY=${ETHOSN_RUNTIME_LIBRARY}") - endif(ETHOSN_FOUND) + endif(NOT ETHOSN_FOUND) endmacro(find_ethosn) From 5e621103571b5fcdc61e195c629d46525d6292e1 Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Wed, 12 Aug 2020 15:26:25 +0100 Subject: [PATCH 30/37] Use GetDataSize Change-Id: I16988f3adbe6e03fc47fa0a77cb5febb7a02eaab --- src/runtime/contrib/ethosn/ethosn_device.cc | 11 ++--------- 1 file changed, 2 insertions(+), 9 deletions(-) diff --git a/src/runtime/contrib/ethosn/ethosn_device.cc b/src/runtime/contrib/ethosn/ethosn_device.cc index 0ec6fa26a3fc..137ab85394e6 100644 --- a/src/runtime/contrib/ethosn/ethosn_device.cc +++ b/src/runtime/contrib/ethosn/ethosn_device.cc @@ -26,6 +26,7 @@ #include #include #include +#include #include #include @@ -45,14 +46,6 @@ namespace ethosn { namespace sl = ::ethosn::support_library; namespace dl = ::ethosn::driver_library; -int64_t GetTensorSize(const DLTensor& tensor) { - int64_t size = 1; - for (int i = 0; i < tensor.ndim; i++) { - size *= tensor.shape[i]; - } - return size; -} - bool WaitForInference(dl::Inference* inference, int timeout) { // Wait for inference to complete int fd = inference->GetFileDescriptor(); @@ -97,7 +90,7 @@ void CreateBuffers(std::vector >* fm, for (auto buffer : tensors) { auto* data = static_cast(buffer->data); // The NPU only needs the size of the tensor * uint8_t. - auto data_size = static_cast(GetTensorSize(*buffer)); + auto data_size = static_cast(GetDataSize(*buffer)); (*fm)[index++] = std::make_shared(data, data_size, dl::DataFormat::NHWC); } } From d9b5e2205b1a3aff1b513dd1683ce4b3cc22febb Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Wed, 12 Aug 2020 15:28:15 +0100 Subject: [PATCH 31/37] Use const& Change-Id: I664982d219f9a7d1f961dbfe84d12f66e2e5f5cb --- src/runtime/contrib/ethosn/ethosn_device.cc | 8 +++++--- src/runtime/contrib/ethosn/ethosn_device.h | 2 +- 2 files changed, 6 insertions(+), 4 deletions(-) diff --git a/src/runtime/contrib/ethosn/ethosn_device.cc b/src/runtime/contrib/ethosn/ethosn_device.cc index 137ab85394e6..7e0d43fcf4cb 100644 --- a/src/runtime/contrib/ethosn/ethosn_device.cc +++ b/src/runtime/contrib/ethosn/ethosn_device.cc @@ -24,9 +24,9 @@ #include #include +#include #include #include -#include #include #include @@ -96,7 +96,8 @@ void CreateBuffers(std::vector >* fm, } bool Inference(tvm::runtime::TVMArgs args, sl::CompiledNetwork* network, - std::vector input_order, std::vector output_order) { + const std::vector& input_order, + const std::vector& output_order) { // Unpack parameters uint8_t argc = 0; std::vector inputs(input_order.size()); @@ -197,7 +198,8 @@ TVM_REGISTER_GLOBAL("relay.ethos-n.test.infra.inference_result") // Allow the ethos-n support code to be tested without a device bool Inference(tvm::runtime::TVMArgs args, sl::CompiledNetwork* network, - std::vector input_order, std::vector output_order) { + const std::vector& input_order, + const std::vector& output_order) { std::vector outputs; for (int argc = network->GetInputBufferInfos().size(); argc < args.size(); argc++) { outputs.push_back(args[argc]); diff --git a/src/runtime/contrib/ethosn/ethosn_device.h b/src/runtime/contrib/ethosn/ethosn_device.h index de62f60fd384..d631d242e993 100644 --- a/src/runtime/contrib/ethosn/ethosn_device.h +++ b/src/runtime/contrib/ethosn/ethosn_device.h @@ -35,7 +35,7 @@ namespace ethosn { namespace sl = ::ethosn::support_library; bool Inference(tvm::runtime::TVMArgs args, sl::CompiledNetwork* network, - std::vector input_order, std::vector output_order); + const std::vector& input_order, const std::vector& output_order); } // namespace ethosn } // namespace runtime From d06da5a84bb57331b95d7b40fa69328791671d91 Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Wed, 12 Aug 2020 17:45:36 +0100 Subject: [PATCH 32/37] Fix python linting Change-Id: Id965ccc037fd40cbdfcb58d922cc8d5fb8c87dfe --- python/tvm/relay/op/contrib/ethosn.py | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/python/tvm/relay/op/contrib/ethosn.py b/python/tvm/relay/op/contrib/ethosn.py index 1b14bdeb6a47..de70297a7889 100644 --- a/python/tvm/relay/op/contrib/ethosn.py +++ b/python/tvm/relay/op/contrib/ethosn.py @@ -16,8 +16,8 @@ # under the License. # pylint: disable=invalid-name, unused-argument """Arm(R) Ethos(TM) -N NPU supported operators.""" -import tvm.ir from enum import Enum +import tvm.ir from ... import qnn as _qnn from . import _ethosn as support @@ -36,9 +36,8 @@ def ethosn_available(): if not tvm.get_global_func("relay.ethos-n.query", True): print("skip because Ethos-N module is not available") return Available.UNAVAILABLE - else: - hw = tvm.get_global_func("relay.ethos-n.query")() - return Available.SW_AND_HW if hw else Available.SW_ONLY + hw = tvm.get_global_func("relay.ethos-n.query")() + return Available.SW_AND_HW if hw else Available.SW_ONLY @tvm.ir.register_op_attr("qnn.concatenate", "target.ethos-n") From c6a0c1b462bdbf9f72a8263a4b3b0dc528cff20f Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Wed, 12 Aug 2020 19:03:31 +0100 Subject: [PATCH 33/37] Remove load/save to file Change-Id: I7f8c3f5c8948c3f15551d28e3fee6e00120663ef --- src/runtime/contrib/ethosn/ethosn_runtime.cc | 19 ------------------- src/runtime/contrib/ethosn/ethosn_runtime.h | 11 ----------- 2 files changed, 30 deletions(-) diff --git a/src/runtime/contrib/ethosn/ethosn_runtime.cc b/src/runtime/contrib/ethosn/ethosn_runtime.cc index 98301f235e58..0fbebcf16139 100644 --- a/src/runtime/contrib/ethosn/ethosn_runtime.cc +++ b/src/runtime/contrib/ethosn/ethosn_runtime.cc @@ -120,25 +120,6 @@ Module EthosnModule::LoadFromBinary(void* strm) { return Module(n); } -void EthosnModule::SaveToFile(const std::string& path, const std::string& format) { - std::string data; - dmlc::MemoryStringStream writer(&data); - dmlc::SeekStream* strm = &writer; - SaveToBinary(strm); - SaveBinaryToFile(path, data); -} - -Module EthosnModule::LoadFromFile(const std::string& path) { - std::string data; - LoadBinaryFromFile(path, &data); - dmlc::MemoryStringStream reader(&data); - return LoadFromBinary(&reader); -} - -TVM_REGISTER_GLOBAL("runtime.module.loadfile_ethos-n").set_body([](TVMArgs args, TVMRetValue* rv) { - *rv = EthosnModule::LoadFromFile(args[0]); -}); - TVM_REGISTER_GLOBAL("runtime.module.loadbinary_ethos-n") .set_body([](TVMArgs args, TVMRetValue* rv) { *rv = EthosnModule::LoadFromBinary(args[0]); }); } // namespace ethosn diff --git a/src/runtime/contrib/ethosn/ethosn_runtime.h b/src/runtime/contrib/ethosn/ethosn_runtime.h index 39cdba8cdacc..730739cbd230 100644 --- a/src/runtime/contrib/ethosn/ethosn_runtime.h +++ b/src/runtime/contrib/ethosn/ethosn_runtime.h @@ -86,17 +86,6 @@ class EthosnModule : public ModuleNode { * ] * number of functions */ static Module LoadFromBinary(void* strm); - /*! - * \brief Save a module to a specified path. - * \param path Where to save the serialized module. - */ - void SaveToFile(const std::string& path, const std::string& format) override; - /*! - * \brief Create a module from a file. - * \param path The path of the file containing the serialized module. - * \return The created Ethos-N module. - */ - static Module LoadFromFile(const std::string& path); const char* type_key() const override { return "ethos-n"; } From 4fe1a6fd1e575eb89e0be51d98c77ceca9989fdd Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Wed, 12 Aug 2020 19:05:45 +0100 Subject: [PATCH 34/37] data->data Change-Id: Ifb861ebbfeaaf4b154f4b1515f83a46aecf86e50 --- src/relay/backend/contrib/ethosn/ethosn_api.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/relay/backend/contrib/ethosn/ethosn_api.cc b/src/relay/backend/contrib/ethosn/ethosn_api.cc index cbd10bb9d743..d92e35afeea0 100644 --- a/src/relay/backend/contrib/ethosn/ethosn_api.cc +++ b/src/relay/backend/contrib/ethosn/ethosn_api.cc @@ -154,7 +154,7 @@ EthosnError EthosnAPI::AsConstant(const Expr& expr, T* out) { return EthosnError("expected constant data"); } runtime::NDArray data = Downcast(expr)->data; - *out = *static_cast(data.operator->()->data); + *out = *static_cast(data->data); return EthosnError(); } From ae17572aa8f2de6a45360d26b369a1535101fc04 Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Wed, 12 Aug 2020 19:07:03 +0100 Subject: [PATCH 35/37] Remove specific cpu target Change-Id: I920568cc7a81cd77d44f8604f571340a330f3e62 --- tests/python/contrib/test_ethosn/infrastructure.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/python/contrib/test_ethosn/infrastructure.py b/tests/python/contrib/test_ethosn/infrastructure.py index 257e16c124b4..ed6631499c11 100644 --- a/tests/python/contrib/test_ethosn/infrastructure.py +++ b/tests/python/contrib/test_ethosn/infrastructure.py @@ -89,7 +89,7 @@ def build(mod, params, npu=True, expected_host_ops=0, npu_partitions=1): with tvm.transform.PassContext(opt_level=3, config={ "relay.ext.ethos-n.options": {"variant": 0} }): - with tvm.target.create("llvm -mcpu=core-avx2"): + with tvm.target.create("llvm"): if npu: f = relay.build_module.bind_params_by_name(mod["main"], params) mod = tvm.IRModule() From 57d1569876fa6926fb7ce8815a63469011638fad Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Thu, 13 Aug 2020 10:00:28 +0100 Subject: [PATCH 36/37] Test export/load module Change-Id: Ib605458127485e2015ac012ec515ced5900705f3 --- tests/python/contrib/test_ethosn/infrastructure.py | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/tests/python/contrib/test_ethosn/infrastructure.py b/tests/python/contrib/test_ethosn/infrastructure.py index ed6631499c11..c6278334cfec 100644 --- a/tests/python/contrib/test_ethosn/infrastructure.py +++ b/tests/python/contrib/test_ethosn/infrastructure.py @@ -112,6 +112,12 @@ def build(mod, params, npu=True, expected_host_ops=0, npu_partitions=1): def run(graph, lib, params, inputs, outputs, npu=True): + # Export and load lib to confirm this works + lib_name = "mod.so" + temp = util.tempdir() + lib_path = temp.relpath(lib_name) + lib.export_library(lib_path) + lib = tvm.runtime.load_module(lib_path) module = graph_runtime.create(graph, lib, tvm.cpu()) module.set_input(**inputs) module.set_input(**params) From 3ab99dde4a73eb3973133bbf95fedbeeb3b82bfa Mon Sep 17 00:00:00 2001 From: Matthew Barrett Date: Thu, 13 Aug 2020 10:07:10 +0100 Subject: [PATCH 37/37] Fix cmake garbage Change-Id: I32f3c967192c7c278ef33c52cac5fb5da682cd1b --- CMakeLists.txt | 4 ---- 1 file changed, 4 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index e4d40a560afd..ce678ea1bc7d 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -309,11 +309,7 @@ include(cmake/modules/Metal.cmake) include(cmake/modules/ROCM.cmake) include(cmake/modules/LLVM.cmake) include(cmake/modules/Micro.cmake) -<<<<<<< HEAD -======= -include(cmake/modules/ANTLR.cmake) include(cmake/modules/contrib/EthosN.cmake) ->>>>>>> [BYOC][ETHOSN] Introduce the Ethos-N BYOC integration include(cmake/modules/contrib/BLAS.cmake) include(cmake/modules/contrib/CODEGENC.cmake) include(cmake/modules/contrib/DNNL.cmake)