Skip to content

Commit 3cff742

Browse files
mbaretleo-blonktristan-armleandronRamana Radhakrishnan
authored
[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 <[email protected]> Co-authored-by: Tristan O'Connor <[email protected]> Co-authored-by: Leandro Nunes <[email protected]> Co-authored-by: Ramana Radhakrishnan <[email protected]> Co-authored-by: Luke Hutton <[email protected]> * 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 <[email protected]> Co-authored-by: Tristan O'Connor <[email protected]> Co-authored-by: Leandro Nunes <[email protected]> Co-authored-by: Ramana Radhakrishnan <[email protected]> Co-authored-by: Luke Hutton <[email protected]>
1 parent 939a42b commit 3cff742

File tree

24 files changed

+2405
-0
lines changed

24 files changed

+2405
-0
lines changed

‎CMakeLists.txt‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -8,6 +8,7 @@ include(cmake/util/FindOpenCL.cmake)
88
include(cmake/util/FindVulkan.cmake)
99
include(cmake/util/FindLLVM.cmake)
1010
include(cmake/util/FindROCM.cmake)
11+
include(cmake/util/FindEthosN.cmake)
1112

1213
if(EXISTS ${CMAKE_CURRENT_BINARY_DIR}/config.cmake)
1314
include(${CMAKE_CURRENT_BINARY_DIR}/config.cmake)
@@ -44,6 +45,7 @@ tvm_option(INSTALL_DEV "Install compiler infrastructure" OFF)
4445
tvm_option(HIDE_PRIVATE_SYMBOLS "Compile with -fvisibility=hidden." OFF)
4546
tvm_option(USE_TF_TVMDSOOP "Build with TensorFlow TVMDSOOp" OFF)
4647
tvm_option(USE_FALLBACK_STL_MAP "Use TVM's POD compatible Map" OFF)
48+
tvm_option(USE_ETHOSN "Build with Arm Ethos-N" OFF)
4749

4850
# 3rdparty libraries
4951
tvm_option(DLPACK_PATH "Path to DLPACK" "3rdparty/dlpack/include")
@@ -307,6 +309,7 @@ include(cmake/modules/Metal.cmake)
307309
include(cmake/modules/ROCM.cmake)
308310
include(cmake/modules/LLVM.cmake)
309311
include(cmake/modules/Micro.cmake)
312+
include(cmake/modules/contrib/EthosN.cmake)
310313
include(cmake/modules/contrib/BLAS.cmake)
311314
include(cmake/modules/contrib/CODEGENC.cmake)
312315
include(cmake/modules/contrib/DNNL.cmake)

‎cmake/config.cmake‎

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -210,6 +210,16 @@ set(USE_DNNL_CODEGEN OFF)
210210
set(USE_ARM_COMPUTE_LIB OFF)
211211
set(USE_ARM_COMPUTE_LIB_GRAPH_RUNTIME OFF)
212212

213+
# Whether to build with Arm Ethos-N support
214+
# Possible values:
215+
# - OFF: disable Arm Ethos-N support
216+
# - path/to/arm-ethos-N-stack: use a specific version of the
217+
# Ethos-N driver stack
218+
set(USE_ETHOSN OFF)
219+
# If USE_ETHOSN is enabled, use ETHOSN_HW (ON) if Ethos-N hardware is available on this machine
220+
# otherwise use ETHOSN_HW (OFF) to use the software test infrastructure
221+
set(USE_ETHOSN_HW OFF)
222+
213223
# Build ANTLR parser for Relay text format
214224
# Possible values:
215225
# - ON: enable ANTLR by searching default locations (cmake find_program for antlr4 and /usr/local for jar)

‎cmake/modules/contrib/EthosN.cmake‎

Lines changed: 57 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,57 @@
1+
# Licensed to the Apache Software Foundation (ASF) under one
2+
# or more contributor license agreements. See the NOTICE file
3+
# distributed with this work for additional information
4+
# regarding copyright ownership. The ASF licenses this file
5+
# to you under the Apache License, Version 2.0 (the
6+
# "License"); you may not use this file except in compliance
7+
# with the License. You may obtain a copy of the License at
8+
#
9+
# http://www.apache.org/licenses/LICENSE-2.0
10+
#
11+
# Unless required by applicable law or agreed to in writing,
12+
# software distributed under the License is distributed on an
13+
# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
14+
# KIND, either express or implied. See the License for the
15+
# specific language governing permissions and limitations
16+
# under the License.
17+
18+
# Arm Ethos-N rules
19+
20+
if(NOT USE_ETHOSN STREQUAL "OFF")
21+
find_ethosn(${USE_ETHOSN})
22+
23+
if(NOT ETHOSN_FOUND)
24+
message(FATAL_ERROR "Cannot find Ethos-N, USE_ETHOSN=" ${USE_ETHOSN})
25+
26+
else()
27+
include_directories(${ETHOSN_INCLUDE_DIRS})
28+
add_definitions(${ETHOSN_DEFINITIONS})
29+
30+
message(STATUS "Build with Ethos-N ${ETHOSN_PACKAGE_VERSION}")
31+
32+
file(GLOB ETHOSN_RUNTIME_CONTRIB_SRC
33+
CONFIGURE_DEPENDS src/runtime/contrib/ethosn/ethosn_runtime.cc
34+
CONFIGURE_DEPENDS src/runtime/contrib/ethosn/ethosn_device.cc)
35+
list(APPEND RUNTIME_SRCS ${ETHOSN_RUNTIME_CONTRIB_SRC})
36+
37+
file(GLOB COMPILER_ETHOSN_SRCS
38+
CONFIGURE_DEPENDS src/relay/backend/contrib/ethosn/*)
39+
list(APPEND COMPILER_SRCS ${COMPILER_ETHOSN_SRCS})
40+
41+
list(APPEND TVM_LINKER_LIBS ${ETHOSN_COMPILER_LIBRARY}
42+
${ETHOSN_RUNTIME_LIBRARY})
43+
list(APPEND TVM_RUNTIME_LINKER_LIBS ${ETHOSN_COMPILER_LIBRARY}
44+
${ETHOSN_RUNTIME_LIBRARY})
45+
46+
if(NOT MSVC)
47+
set_source_files_properties(${COMPILER_ETHOSN_SRCS}
48+
PROPERTIES COMPILE_DEFINITIONS "DMLC_ENABLE_RTTI=0")
49+
set_source_files_properties(${COMPILER_ETHOSN_SRCS}
50+
PROPERTIES COMPILE_FLAGS "-fno-rtti")
51+
endif()
52+
endif(NOT ETHOSN_FOUND)
53+
else()
54+
if(USE_ETHOSN_HW)
55+
message(FATAL_ERROR "Cannot enable Ethos-N HW if USE_ETHOSN=OFF")
56+
endif()
57+
endif(NOT USE_ETHOSN STREQUAL "OFF")

‎cmake/util/FindEthosN.cmake‎

Lines changed: 94 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,94 @@
1+
# Licensed to the Apache Software Foundation (ASF) under one
2+
# or more contributor license agreements. See the NOTICE file
3+
# distributed with this work for additional information
4+
# regarding copyright ownership. The ASF licenses this file
5+
# to you under the Apache License, Version 2.0 (the
6+
# "License"); you may not use this file except in compliance
7+
# with the License. You may obtain a copy of the License at
8+
#
9+
# http://www.apache.org/licenses/LICENSE-2.0
10+
#
11+
# Unless required by applicable law or agreed to in writing,
12+
# software distributed under the License is distributed on an
13+
# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
14+
# KIND, either express or implied. See the License for the
15+
# specific language governing permissions and limitations
16+
# under the License.
17+
18+
#######################################################
19+
# Find Arm Ethos-N libraries
20+
#
21+
# Usage:
22+
# find_ethosn(${USE_ETHOSN})
23+
#
24+
# - When USE_ETHOSN=/path/to/ethos-sdk-path, use the path from USE_ETHOSN
25+
# - Else, when environment variable ETHOSN_STACK is set, use that path
26+
# - When USE_ETHOSN=ON, use auto search
27+
#
28+
# Provide variables:
29+
#
30+
# - ETHOSN_FOUND
31+
# - ETHOSN_PACKAGE_VERSION
32+
# - ETHOSN_DEFINITIONS
33+
# - ETHOSN_INCLUDE_DIRS
34+
# - ETHOSN_COMPILER_LIBRARY
35+
# - ETHOSN_RUNTIME_LIBRARY
36+
37+
macro(find_ethosn use_ethosn)
38+
set(__use_ethosn ${use_ethosn})
39+
if(IS_DIRECTORY ${__use_ethosn})
40+
set(__ethosn_stack ${__use_ethosn})
41+
message(STATUS "Arm Ethos-N driver stack PATH=" ${__use_ethosn})
42+
elseif(IS_DIRECTORY $ENV{ETHOSN_STACK})
43+
set(__ethosn_stack $ENV{ETHOSN_STACK})
44+
message(STATUS "Arm Ethos-N driver stack from env=" ${__use_ethosn})
45+
else()
46+
set(__ethosn_stack "")
47+
endif()
48+
49+
if(__ethosn_stack)
50+
set(ETHOSN_INCLUDE_DIRS "")
51+
# Compile-time support
52+
find_path(_SL_DIR NAMES Support.hpp
53+
PATHS ${__ethosn_stack}/include/ethosn_support_library)
54+
string(REGEX REPLACE "/ethosn_support_library" "" _SL_DIR2 ${_SL_DIR})
55+
list(APPEND ETHOSN_INCLUDE_DIRS "${_SL_DIR2}")
56+
57+
find_library(ETHOSN_COMPILER_LIBRARY NAMES EthosNSupport
58+
PATHS ${__ethosn_stack}/lib)
59+
find_library(ETHOSN_COMPILER_LIBRARY NAMES EthosNSupport)
60+
61+
set(ETHOSN_PACKAGE_VERSION "0.1.1")
62+
63+
if(USE_ETHOSN_HW STREQUAL "ON")
64+
# Runtime hardware support
65+
find_path(_DL_DIR NAMES Network.hpp
66+
PATHS ${__ethosn_stack}/include/ethosn_driver_library)
67+
string(REGEX REPLACE "/ethosn_driver_library" "" _DL_DIR2 ${_DL_DIR})
68+
list(APPEND ETHOSN_INCLUDE_DIRS "${_DL_DIR2}")
69+
70+
find_library(ETHOSN_RUNTIME_LIBRARY NAMES EthosNDriver
71+
PATHS ${__ethosn_stack}/lib)
72+
find_library(ETHOSN_RUNTIME_LIBRARY NAMES EthosNDriver)
73+
set(ETHOSN_DEFINITIONS -DETHOSN_HW)
74+
endif ()
75+
76+
if(ETHOSN_COMPILER_LIBRARY)
77+
set(ETHOSN_FOUND TRUE)
78+
endif()
79+
endif(__ethosn_stack)
80+
81+
if(NOT ETHOSN_FOUND)
82+
if(__use_ethosn STREQUAL "ON")
83+
message(WARNING "No cmake find_package available for Arm Ethos-N")
84+
endif()
85+
86+
# additional libraries
87+
else()
88+
message(STATUS "Found ETHOSN_DEFINITIONS=${ETHOSN_DEFINITIONS}")
89+
message(STATUS "Found ETHOSN_INCLUDE_DIRS=${ETHOSN_INCLUDE_DIRS}")
90+
message(STATUS "Found ETHOSN_COMPILER_LIBRARY=${ETHOSN_COMPILER_LIBRARY}")
91+
message(STATUS "Found ETHOSN_RUNTIME_LIBRARY=${ETHOSN_RUNTIME_LIBRARY}")
92+
endif(NOT ETHOSN_FOUND)
93+
94+
endmacro(find_ethosn)

‎python/tvm/relay/op/contrib/__init__.py‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -21,3 +21,4 @@
2121
from .arm_compute_lib import *
2222
from .dnnl import *
2323
from .coreml import *
24+
from .ethosn import *

‎python/tvm/relay/op/contrib/_ethosn.py‎

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
# Licensed to the Apache Software Foundation (ASF) under one
2+
# or more contributor license agreements. See the NOTICE file
3+
# distributed with this work for additional information
4+
# regarding copyright ownership. The ASF licenses this file
5+
# to you under the Apache License, Version 2.0 (the
6+
# "License"); you may not use this file except in compliance
7+
# with the License. You may obtain a copy of the License at
8+
#
9+
# http://www.apache.org/licenses/LICENSE-2.0
10+
#
11+
# Unless required by applicable law or agreed to in writing,
12+
# software distributed under the License is distributed on an
13+
# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
14+
# KIND, either express or implied. See the License for the
15+
# specific language governing permissions and limitations
16+
# under the License.
17+
18+
"""Expose 'is supported' functions to Python."""
19+
20+
import tvm._ffi
21+
22+
tvm._ffi._init_api("relay.ethos-n.support", __name__)

‎python/tvm/relay/op/contrib/ethosn.py‎

Lines changed: 89 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,89 @@
1+
# Licensed to the Apache Software Foundation (ASF) under one
2+
# or more contributor license agreements. See the NOTICE file
3+
# distributed with this work for additional information
4+
# regarding copyright ownership. The ASF licenses this file
5+
# to you under the Apache License, Version 2.0 (the
6+
# "License"); you may not use this file except in compliance
7+
# with the License. You may obtain a copy of the License at
8+
#
9+
# http://www.apache.org/licenses/LICENSE-2.0
10+
#
11+
# Unless required by applicable law or agreed to in writing,
12+
# software distributed under the License is distributed on an
13+
# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
14+
# KIND, either express or implied. See the License for the
15+
# specific language governing permissions and limitations
16+
# under the License.
17+
# pylint: disable=invalid-name, unused-argument
18+
"""Arm(R) Ethos(TM) -N NPU supported operators."""
19+
from enum import Enum
20+
import tvm.ir
21+
from ... import qnn as _qnn
22+
from . import _ethosn as support
23+
24+
25+
class Available(Enum):
26+
UNAVAILABLE = 0
27+
SW_ONLY = 1
28+
SW_AND_HW = 2
29+
30+
def __bool__(self):
31+
return self != Available.UNAVAILABLE
32+
33+
34+
def ethosn_available():
35+
"""Return whether Ethos-N software and hardware support is available"""
36+
if not tvm.get_global_func("relay.ethos-n.query", True):
37+
print("skip because Ethos-N module is not available")
38+
return Available.UNAVAILABLE
39+
hw = tvm.get_global_func("relay.ethos-n.query")()
40+
return Available.SW_AND_HW if hw else Available.SW_ONLY
41+
42+
43+
@tvm.ir.register_op_attr("qnn.concatenate", "target.ethos-n")
44+
def qnn_concatenate(attrs, args):
45+
"""Check if a concatenate is supported by Ethos-N."""
46+
if not ethosn_available():
47+
return False
48+
49+
conc = _qnn.op.concatenate(*args, **attrs)
50+
if not support.concatenate(conc):
51+
return False
52+
53+
# Support library has some unenforced restrictions on qnn params
54+
min_range = 1e9
55+
max_range = -1e9
56+
qnn_params = []
57+
for i in range(len(args[1].fields)):
58+
scale = args[1].fields[i].data.asnumpy()
59+
zero_point = args[2].fields[i].data.asnumpy()
60+
min_range = min(-1 * zero_point * scale, min_range)
61+
max_range = max((255 - zero_point) * scale, max_range)
62+
qnn_params.append((scale, zero_point))
63+
64+
scale = (max_range - min_range) / 255
65+
zero_point = int(-min_range/scale)
66+
if (scale, zero_point) in qnn_params:
67+
return True
68+
69+
return False
70+
71+
72+
@tvm.ir.register_op_attr("split", "target.ethos-n")
73+
def split(attrs, args):
74+
"""Check if a split is supported by Ethos-N."""
75+
if not ethosn_available():
76+
return False
77+
78+
if isinstance(attrs["indices_or_sections"], tvm.tir.IntImm):
79+
sp = tvm.relay.split(*args,
80+
indices_or_sections=attrs["indices_or_sections"].value,
81+
axis=attrs["axis"])
82+
else:
83+
sp = tvm.relay.split(*args,
84+
indices_or_sections=attrs["indices_or_sections"],
85+
axis=attrs["axis"])
86+
if not support.split(sp.astuple()):
87+
return False
88+
89+
return True

‎src/relay/backend/contrib/ethosn/capabilities.h‎

Lines changed: 81 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,81 @@
1+
/*
2+
* Licensed to the Apache Software Foundation (ASF) under one
3+
* or more contributor license agreements. See the NOTICE file
4+
* distributed with this work for additional information
5+
* regarding copyright ownership. The ASF licenses this file
6+
* to you under the Apache License, Version 2.0 (the
7+
* "License"); you may not use this file except in compliance
8+
* with the License. You may obtain a copy of the License at
9+
*
10+
* http://www.apache.org/licenses/LICENSE-2.0
11+
*
12+
* Unless required by applicable law or agreed to in writing,
13+
* software distributed under the License is distributed on an
14+
* "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
15+
* KIND, either express or implied. See the License for the
16+
* specific language governing permissions and limitations
17+
* under the License.
18+
*/
19+
20+
/*!
21+
* \file src/relay/backend/contrib/ethosn/capabilities.h
22+
* \brief The Ethos-N processor series has four variants, the Ethos-N37, Ethos-N57, Ethos-N77
23+
* and the Ethos-N78. This release of the integration supports the first three variants.
24+
* Configuration information for each variant is stored as a blob in this file. These blobs
25+
* are passed into the Ethos-N support library, which in turn uses them to optimize the
26+
* generated command-stream appropriately for the specified variant.
27+
*/
28+
29+
#ifndef TVM_RELAY_BACKEND_CONTRIB_ETHOSN_CAPABILITIES_H_
30+
#define TVM_RELAY_BACKEND_CONTRIB_ETHOSN_CAPABILITIES_H_
31+
32+
#include <vector>
33+
34+
namespace tvm {
35+
namespace relay {
36+
namespace contrib {
37+
namespace ethosn {
38+
39+
/* Ethos-N variants (N77, N57 and N37)
40+
* variant[0] - N77
41+
* variant[1] - N57
42+
* variant[2] - N37
43+
*/
44+
static std::vector<char> variants[3] = {
45+
{
46+
0x02, 0x00, 0x00, 0x00, 0x74, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
47+
0x00, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x10, 0x00, 0x10, 0x00,
48+
0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x00,
49+
0x10, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00,
50+
0x01, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x10, 0x00, 0x00,
51+
0x00, 0x01, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x01, 0x00,
52+
0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x40, 0x00, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x01,
53+
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
54+
},
55+
{
56+
0x02, 0x00, 0x00, 0x00, 0x74, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
57+
0x00, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x10, 0x00, 0x10, 0x00,
58+
0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x00,
59+
0x10, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00,
60+
0x01, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x10, 0x00, 0x00,
61+
0x00, 0x01, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x01, 0x00,
62+
0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x40, 0x00, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x01,
63+
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
64+
},
65+
{
66+
0x02, 0x00, 0x00, 0x00, 0x74, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
67+
0x00, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x10, 0x00, 0x10, 0x00,
68+
0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x01, 0x00, 0x00, 0x00, 0x00,
69+
0x10, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00,
70+
0x01, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x10, 0x00, 0x00,
71+
0x00, 0x01, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x04, 0x00, 0x00, 0x00, 0x01, 0x00,
72+
0x00, 0x00, 0x08, 0x00, 0x00, 0x00, 0x40, 0x00, 0x00, 0x00, 0x00, 0x02, 0x00, 0x00, 0x01,
73+
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
74+
}};
75+
76+
} // namespace ethosn
77+
} // namespace contrib
78+
} // namespace relay
79+
} // namespace tvm
80+
81+
#endif // TVM_RELAY_BACKEND_CONTRIB_ETHOSN_CAPABILITIES_H_

0 commit comments

Comments
 (0)