[BYOC][ETHOSN] Introduce the Ethos-N BYOC integration (apache#6222)

* [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 <Leo.Blonk@arm.com> Co-authored-by: Tristan O'Connor <tristan.oconnor@arm.com> Co-authored-by: Leandro Nunes <leandro.nunes@arm.com> Co-authored-by: Ramana Radhakrishnan <ramana.radhakrishnan@arm.com> Co-authored-by: Luke Hutton <luke.hutton@arm.com> * Turn off USE_ETHOSN_HW by default Change-Id: Ie2ce4528e16e93aa83df46f8a229c0ce89b45252 * Update capabilities file Change-Id: Iebd0c62d6bc7e446662abdee4882ac874ad98aa3 * Fix missing header Change-Id: I0c89e380dd1d795755a1884c06a7b317a99fe297 * Update cmake comments on ETHOSN_HW Change-Id: I2e96a1c818a82e5174fd94e483b0bdb3e4375a7d * Add checker for case when USE_ETHOSN=OFF and USE_ETHOSN_HW=ON Change-Id: Id5c9cfb866914a0298b44ead40fcbe3764ce443c * Fix 'available' boolean Change-Id: I78e54fb9f472d2815886bea4d94b7247e0d129de * Check availability in op registration Change-Id: Iecfea7dca7301dd684199c9b32f99f2113fdfd56 * Remove unnecessary line Change-Id: Idf5cab853027adb0b0292de877e6dc02683821d7 * Simplify getting output_size Change-Id: If4643924768c2d7ea98525e9f792b7223cc2bcdf * Remove unnecessary new line Change-Id: Ia689c59cac28bd91e237ceecd829d8cf56d0d9c1 * Remove NOLINTS Change-Id: I149b97b28b516c7d9288a0858b2fbf1497e70250 * Remove unused parts of PR Change-Id: I2db5b89d8fe2c114ab92305cdcf06d0fc45f4d2a * Fix CI Ethos-N settings Change-Id: Idd955755d6f6d1cd3843462f627d0d952729e467 * Removed unnecessary line in infra Change-Id: I0ea866adf5d9166db85dd82d013a631d991ae633 * Remove unnecessary len in infra Change-Id: I869e8233d41c6ab7c2dc80f47d976c974043b80c * Rename 'cpu_ops' to 'host_ops' Change-Id: I79a6ffcfd48cd055d279f493c672ec82f0c68e5c * Added explanation on mocking Change-Id: I1e88c07a47464e44cb45c6a327ec9c7e2d70cc94 * IsEthosOp -> IsEthosnOp Change-Id: I4fc1b462a74f8fae231ebafac614dd8d45be0feb * Improve documentation in ethosn_api.h Change-Id: I5586a7ba7ce71da667a6a9c6dd2e591028eb43b2 * No longer iterate over module when compiling Change-Id: I80e1d494c6d574be06a2375e831343485712914d * Move EthosnCompiler implementations into codegen.cc Change-Id: I5bb6e9f62722d930d9dc040ac62bf87f29dd74c5 * Fix linting Change-Id: Ia44ec741a5330ad289cc6b5cd2bb1ed784fe6afc * Refactor EthosnAPI compilation functions into EthosnCompiler Change-Id: Iee0aecbe43a84fefb437ab9ff064e3f8b42c80a4 * Improve docs for Tvm2Npu Change-Id: Ia39e9e1508513ca39c1d585fbccc3ae38fcbb9fb * Move more implementation out of headers Change-Id: I1e33084ceb520b75f06b4d7a4acff5b9b2225bd5 * Move implementation in ethosn_api.h Change-Id: I51ab386892a2aa84aa47d03641aac8468f5737ae * Improve docs for capabilities.h Change-Id: Iaaee508aafa1cbb7650a04ed87bd6c1b91823a58 * Use else() in cmake Change-Id: I4b64a87f32b3616ec87c9937d9fc998b8dc5d7b4 * Use GetDataSize Change-Id: I16988f3adbe6e03fc47fa0a77cb5febb7a02eaab * Use const& Change-Id: I664982d219f9a7d1f961dbfe84d12f66e2e5f5cb * Fix python linting Change-Id: Id965ccc037fd40cbdfcb58d922cc8d5fb8c87dfe * Remove load/save to file Change-Id: I7f8c3f5c8948c3f15551d28e3fee6e00120663ef * data->data Change-Id: Ifb861ebbfeaaf4b154f4b1515f83a46aecf86e50 * Remove specific cpu target Change-Id: I920568cc7a81cd77d44f8604f571340a330f3e62 * Test export/load module Change-Id: Ib605458127485e2015ac012ec515ced5900705f3 * Fix cmake garbage Change-Id: I32f3c967192c7c278ef33c52cac5fb5da682cd1b Co-authored-by: Leo Blonk <Leo.Blonk@arm.com> Co-authored-by: Tristan O'Connor <tristan.oconnor@arm.com> Co-authored-by: Leandro Nunes <leandro.nunes@arm.com> Co-authored-by: Ramana Radhakrishnan <ramana.radhakrishnan@arm.com> Co-authored-by: Luke Hutton <luke.hutton@arm.com>
trevor-m · Aug 26, 2020 · 2146c30 · 2146c30
1 parent aa9cc61
commit 2146c30
Show file tree

Hide file tree

Showing 24 changed files with 2,405 additions and 0 deletions.
diff --git 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)
@@ -45,6 +46,7 @@ tvm_option(HIDE_PRIVATE_SYMBOLS "Compile with -fvisibility=hidden." OFF)
 tvm_option(USE_TF_COMPILE_FLAGS "Build with TensorFlow's compile flags." 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")
@@ -322,6 +324,7 @@ include(cmake/modules/Metal.cmake)
 include(cmake/modules/ROCM.cmake)
 include(cmake/modules/LLVM.cmake)
 include(cmake/modules/Micro.cmake)
+include(cmake/modules/contrib/EthosN.cmake)
 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
@@ -214,6 +214,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 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
 # 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
@@ -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.
+
+# 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})
+
+  else()
+    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(NOT 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")
diff --git a/cmake/util/FindEthosN.cmake b/cmake/util/FindEthosN.cmake
@@ -0,0 +1,94 @@
+# 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()
+
+  # additional libraries
+  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(NOT ETHOSN_FOUND)
+
+endmacro(find_ethosn)
diff --git 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
@@ -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
@@ -0,0 +1,89 @@
+# 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."""
+from enum import Enum
+import tvm.ir
+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
+    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
+
+    # 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 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,
+                             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
@@ -0,0 +1,81 @@
+/*
+ * 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/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_
+
+#include <vector>
+
+namespace tvm {
+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<char> 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,
+        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,
+    }};
+
+}  // namespace ethosn
+}  // namespace contrib
+}  // namespace relay
+}  // namespace tvm
+
+#endif  // TVM_RELAY_BACKEND_CONTRIB_ETHOSN_CAPABILITIES_H_