You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by GitBox <gi...@apache.org> on 2021/11/05 11:00:22 UTC
[GitHub] [tvm] ekalda opened a new pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
ekalda opened a new pull request #9457:
URL: https://github.com/apache/tvm/pull/9457
* Add the ethosu.identity operator which returns the input tensor
* Add an opportunity to requantize the tensor
* Add legalization for reshape and strided slice
* Add a pass that puts an indentity op after a no-op
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] lhutton1 commented on a change in pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
Posted by GitBox <gi...@apache.org>.
lhutton1 commented on a change in pull request #9457:
URL: https://github.com/apache/tvm/pull/9457#discussion_r744641200
##########
File path: tests/python/contrib/test_ethosu/test_type_inference.py
##########
@@ -14,6 +14,14 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
+from tvm import relay, TVMError
+from tvm import relay
+import tvm
+from tvm.relay.testing import run_opt_pass
+from .infra import make_ethosu_conv2d
+from .infra import make_ethosu_pooling
+from .infra import make_ethosu_identity
+
Review comment:
Looks like everything got imported twice, probably when rebasing?
##########
File path: python/tvm/relay/backend/contrib/ethosu/legalize.py
##########
@@ -123,6 +123,109 @@ def __call__(self, *args, **kwargs):
pass
+class StridedSliceRewriter(DFPatternCallback):
+ """This pass brings the strided slice out of the partitioned function"""
+
+ def __init__(self):
+ super().__init__(require_type=True, rewrite_once=True)
+ self.pattern = (wildcard().has_attr({"Composite": "ethosu.strided_slice"}))(wildcard())
+
+ def callback(
+ self, pre: tvm.relay.Expr, post: tvm.relay.Expr, node_map: tvm.ir.container.Map
+ ) -> tvm.relay.Expr:
+ slice_input = post.args[0]
+ attrs = post.op.body.attrs
+ begin = attrs.begin
+ end = attrs.end
+ strides = attrs.strides
+ axes = attrs.axes
+ slice_mode = attrs.slice_mode
+ strided_slice = relay.op.strided_slice(
+ slice_input, begin, end, strides=strides, axes=axes, slice_mode=slice_mode
+ )
+ return strided_slice
+
+
+@ir.transform.module_pass(opt_level=1)
+class LegalizeStridedSlice:
+ """This is the pass that wraps StridedSliceRewriter"""
+
+ def transform_module(
+ self, mod: tvm.ir.IRModule, ctx: tvm.ir.transform.PassContext
+ ) -> tvm.ir.IRModule:
+ for global_var, func in mod.functions.items():
+ func = rewrite(StridedSliceRewriter(), func)
+ mod.update_func(global_var, func)
+ return mod
+
+ def __call__(self, *args, **kwargs):
+ pass
+
+
+class ReshapeRewriter(DFPatternCallback):
+ """This pass brings the reshape out of the partitioned function"""
+
+ def __init__(self):
+ super().__init__(require_type=True, rewrite_once=True)
+ self.pattern = (wildcard().has_attr({"Composite": "ethosu.reshape"}))(wildcard())
+
+ def callback(
+ self, pre: tvm.relay.Expr, post: tvm.relay.Expr, node_map: tvm.ir.container.Map
+ ) -> tvm.relay.Expr:
+ reshape_input = post.args[0]
+ new_shape = post.op.body.attrs.newshape
+ reshape = relay.op.reshape(reshape_input, newshape=new_shape)
+ return reshape
+
+
+@ir.transform.module_pass(opt_level=1)
+class LegalizeReshape:
+ """This is the pass that wraps ReshapeRewriter"""
+
+ def transform_module(
+ self, mod: tvm.ir.IRModule, ctx: tvm.ir.transform.PassContext
+ ) -> tvm.ir.IRModule:
+ for global_var, func in mod.functions.items():
+ func = rewrite(ReshapeRewriter(), func)
+ mod.update_func(global_var, func)
+ return mod
+
+ def __call__(self, *args, **kwargs):
+ pass
+
+
+class NoOpRewriter(DFPatternCallback):
+ """This pass adds and idenity operator to reshape and strided slice to avoid a no op
Review comment:
```suggestion
"""This pass adds an identity operator to reshape and strided slice to avoid a no op
```
##########
File path: python/tvm/relay/backend/contrib/ethosu/te/identity.py
##########
@@ -0,0 +1,73 @@
+# 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
+"""Tensor Expression for identity"""
+from tvm import te
+from .dma import read_compute, write_compute
+
+
+def identity_compute(
+ ifm: te.Tensor,
+ lut: te.Tensor,
+ ifm_scale: float,
+ ifm_zero_point: int,
+ ofm_scale: float,
+ ofm_zero_point: int,
+ activation: str,
+) -> te.Tensor:
+ """A compute operator for the NPU identity operator.
+
+ Parameters
+ ----------
+ ifm : te.Tensor
+ The Input Feature Map tensor (IFM).
+ lut : te.Tensor
+ The look-up table values to use if activation is "LUT", "TANH" or "SIGMOID".
+ ifm_scale : float
+ The quantization scale for the Input Feature Map tensor.
+ ifm_zero_point : int
+ The quantization zero point for the Input Feature Map tensor.
+ ofm_scale : float
+ The quantization scale for the Output Feature Map tensor.
+ ofm_zero_point : int
+ The quantization zero point for the Output Feature Map tensor.
+ activation : str
+ The activation function to use.
+ "NONE" - no activation function.
+ "TANH" - tanh activation function.
+ "SIGMOID" - sigmoid activation function.
+ "LUT" - use a look-up table to perform the activation function.
+
+ Returns
+ -------
+ te.Tensor
+ The OFM tensor.
Review comment:
Nit: Output Feature Map is probably better here since OFM is not yet defined.
##########
File path: src/relay/op/contrib/ethosu/identity.cc
##########
@@ -0,0 +1,128 @@
+/*
+ * 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/op/contrib/ethosu/identity.cc
+ * \brief Property def of the Arm Ethos-U identity op.
+ */
+#include <tvm/relay/op.h>
+
+#include "common.h"
+
+namespace tvm {
+namespace relay {
+namespace op {
+namespace contrib {
+namespace ethosu {
+
+/*! \brief Attributes used by the Ethos(TM)-U NPU identity operator */
+struct EthosuIdentityAttrs : public tvm::AttrsNode<EthosuIdentityAttrs> {
+ double ifm_scale;
+ int ifm_zero_point;
+ double ofm_scale;
+ int ofm_zero_point;
+ String activation;
+
+ TVM_DECLARE_ATTRS(EthosuIdentityAttrs, "relay.attrs.EthosuIdentityAttrs") {
+ TVM_ATTR_FIELD(ifm_scale).describe("The quantization scale for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ifm_zero_point)
+ .describe("The quantization zero point for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_scale).describe("The quantization scale for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_zero_point)
+ .describe("The quantization zero point for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(activation)
+ .describe(
+ "The activation function to use. "
+ "'NONE' - no activation function. "
+ "'TANH' - tanh activation function. "
+ "'SIGMOID' - sigmoid activation function. "
+ "'LUT' - use a look-up table to perform the activation function.")
+ .set_default("NONE");
+ }
+};
+
+TVM_REGISTER_NODE_TYPE(EthosuIdentityAttrs);
+
+bool EthosuIdentityRel(const Array<Type>& types, int num_inputs, const Attrs& attrs,
+ const TypeReporter& reporter) {
+ int ifm_index = 0;
+ int result_index = 2;
+ ICHECK_EQ(types.size(), result_index + 1);
+
+ const auto* ifm = types[ifm_index].as<TensorTypeNode>();
+ if (ifm == nullptr) return false;
+
+ const auto* param = attrs.as<EthosuIdentityAttrs>();
+
+ ICHECK(param != nullptr) << "EthosuIdentityAttrs cannot be nullptr.";
+
+ if (ifm->dtype != DataType::UInt(8) && ifm->dtype != DataType::Int(8)) {
+ reporter->GetDiagCtx().EmitFatal(
+ Diagnostic::Error(reporter->GetSpan())
+ << "Invalid operator: Expected type(uint8) or type(int8) for ifm but was " << ifm->dtype);
+ return false;
+ }
+
+ if (ifm->shape.size() > 4) {
+ reporter->GetDiagCtx().EmitFatal(
+ Diagnostic::Error(reporter->GetSpan())
+ << "Invalid operator: Input Feature Map should be at most 4 dimensional, but was "
+ << ifm->shape);
+ return false;
+ }
+
+ // Assign ofm type
+ auto ofm_shape = ifm->shape;
+ reporter->Assign(types[result_index], TensorType(ofm_shape, ifm->dtype));
+ return true;
+}
+
+Expr MakeEthosuIdentity(Expr ifm, Expr lut, double ifm_scale, int ifm_zero_point, double ofm_scale,
+ int ofm_zero_point, String activation) {
+ auto attrs = make_object<EthosuIdentityAttrs>();
+ attrs->ifm_scale = ifm_scale;
+ attrs->ifm_zero_point = ifm_zero_point;
+ attrs->ofm_scale = ofm_scale;
+ attrs->ofm_zero_point = ofm_zero_point;
+ attrs->activation = std::move(activation);
+ static const Op& op = Op::Get("contrib.ethosu.identity");
+ return Call(op, {ifm, lut}, Attrs(attrs), {});
+}
+
+TVM_REGISTER_GLOBAL("relay.op._make.ethosu_identity").set_body_typed(MakeEthosuIdentity);
+
+RELAY_REGISTER_OP("contrib.ethosu.identity")
+ .describe(R"code(Identity operator for Ethos-U NPUs.
Review comment:
```suggestion
.describe(R"code(Identity operator for Ethos(TM)-U NPUs.
```
##########
File path: src/relay/op/contrib/ethosu/identity.cc
##########
@@ -0,0 +1,128 @@
+/*
+ * 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/op/contrib/ethosu/identity.cc
+ * \brief Property def of the Arm Ethos-U identity op.
+ */
+#include <tvm/relay/op.h>
+
+#include "common.h"
+
+namespace tvm {
+namespace relay {
+namespace op {
+namespace contrib {
+namespace ethosu {
+
+/*! \brief Attributes used by the Ethos(TM)-U NPU identity operator */
+struct EthosuIdentityAttrs : public tvm::AttrsNode<EthosuIdentityAttrs> {
+ double ifm_scale;
+ int ifm_zero_point;
+ double ofm_scale;
+ int ofm_zero_point;
+ String activation;
+
+ TVM_DECLARE_ATTRS(EthosuIdentityAttrs, "relay.attrs.EthosuIdentityAttrs") {
+ TVM_ATTR_FIELD(ifm_scale).describe("The quantization scale for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ifm_zero_point)
+ .describe("The quantization zero point for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_scale).describe("The quantization scale for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_zero_point)
+ .describe("The quantization zero point for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(activation)
+ .describe(
+ "The activation function to use. "
+ "'NONE' - no activation function. "
+ "'TANH' - tanh activation function. "
+ "'SIGMOID' - sigmoid activation function. "
+ "'LUT' - use a look-up table to perform the activation function.")
+ .set_default("NONE");
+ }
+};
+
+TVM_REGISTER_NODE_TYPE(EthosuIdentityAttrs);
+
+bool EthosuIdentityRel(const Array<Type>& types, int num_inputs, const Attrs& attrs,
+ const TypeReporter& reporter) {
+ int ifm_index = 0;
+ int result_index = 2;
Review comment:
Better to make these `const`
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ekalda commented on a change in pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
Posted by GitBox <gi...@apache.org>.
ekalda commented on a change in pull request #9457:
URL: https://github.com/apache/tvm/pull/9457#discussion_r750118419
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir/identity.py
##########
@@ -0,0 +1,87 @@
+# 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
+"""Extract information from the identity operator in TIR."""
+from typing import Dict, Tuple
+import tvm
+from .dma import get_read_params, get_write_params
+from .spec import SerialKernel, SerialActivation, SerialPooling, SerialPadding
+from .utils import get_op_attrs
+
+
+def get_identity_params(
+ stmt: tvm.tir.AttrStmt,
+ producers: Dict[tvm.tir.Var, tvm.tir.AttrStmt],
+ consumers: Dict[tvm.tir.Var, tvm.tir.AttrStmt],
+) -> Tuple[SerialPooling, tvm.tir.Var, tvm.tir.Var]:
+ """Get the parameters necessary to construct a call_extern for a pooling.
Review comment:
Done
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ekalda commented on a change in pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
Posted by GitBox <gi...@apache.org>.
ekalda commented on a change in pull request #9457:
URL: https://github.com/apache/tvm/pull/9457#discussion_r750118571
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir/identity.py
##########
@@ -0,0 +1,87 @@
+# 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
+"""Extract information from the identity operator in TIR."""
+from typing import Dict, Tuple
+import tvm
+from .dma import get_read_params, get_write_params
+from .spec import SerialKernel, SerialActivation, SerialPooling, SerialPadding
+from .utils import get_op_attrs
+
+
+def get_identity_params(
+ stmt: tvm.tir.AttrStmt,
+ producers: Dict[tvm.tir.Var, tvm.tir.AttrStmt],
+ consumers: Dict[tvm.tir.Var, tvm.tir.AttrStmt],
+) -> Tuple[SerialPooling, tvm.tir.Var, tvm.tir.Var]:
+ """Get the parameters necessary to construct a call_extern for a pooling.
+
+ Parameters
+ ----------
+ stmt : tvm.tir.AttrStmt
+ The outermost attribute statement of a convolution loop nest.
+ producers : Dict[tvm.tir.Var, tvm.tir.AttrStmt]
+ A dictionary to associate pointers with the loop nest
+ that produces their values.
+ consumers : Dict[tvm.tir.Var, tvm.tir.AttrStmt]
+ A dictionary to associate pointers with the loop nest
+ that consumes their values.
+
+ Returns
+ -------
+ SerialPooling
+ The parameters needed to construct a 2D pooling.
+ output_pointer : tvm.tir.Var
+ The output pointer of the pooling operation.
+ replace_pointer : tvm.tir.Var
+ The output pointer of the DMA write operation, which is to replace
+ the pooling output pointer.
+
+ """
+ attrs, _ = get_op_attrs(stmt)
+ # Find the inner loop
+ while hasattr(stmt, "body"):
+ stmt = stmt.body
+
+ input_pointer = stmt.value.buffer_var
+ output_pointer = stmt.buffer_var
+
+ read = producers[input_pointer]
+ write = consumers[output_pointer]
+
+ serial_ifm, _, _ = get_read_params(read)
+ serial_ofm, _, write_output_pointer = get_write_params(write)
+
+ replace_pointer = write_output_pointer
+
+ # TODO (maybe): Support stand alone RELU through clamping in identity
Review comment:
Done
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ekalda commented on a change in pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
Posted by GitBox <gi...@apache.org>.
ekalda commented on a change in pull request #9457:
URL: https://github.com/apache/tvm/pull/9457#discussion_r750116578
##########
File path: python/tvm/relay/backend/contrib/ethosu/legalize.py
##########
@@ -423,11 +526,15 @@ class LegalizeEthosU:
def transform_module(
self, mod: tvm.ir.IRModule, ctx: tvm.ir.transform.PassContext
) -> tvm.ir.IRModule:
+ """Legalize the oerators that can be offloaded to the NPU"""
Review comment:
A different docstring appeared in the rebase, so I kept that that different docstring
##########
File path: python/tvm/relay/op/contrib/ethosu.py
##########
@@ -481,6 +493,8 @@ def pattern_table() -> List[Tuple[str, tvm.relay.dataflow_pattern.DFPattern, Cal
qnn_avgpool2d_pattern(),
lambda pat: AvgPool2DParams(pat).is_valid(),
),
+ ("ethosu.strided_slice", strided_slice_pattern(), lambda pat: True),
Review comment:
Done
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] NicolaLancellotti commented on a change in pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
Posted by GitBox <gi...@apache.org>.
NicolaLancellotti commented on a change in pull request #9457:
URL: https://github.com/apache/tvm/pull/9457#discussion_r743677908
##########
File path: python/tvm/relay/backend/contrib/ethosu/op/identity.py
##########
@@ -0,0 +1,98 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+# pylint: disable=unused-argument
+"""Relay identity operator"""
Review comment:
```suggestion
"""Relay operator for identity for Arm(R) Ethos(TM)-U NPU"""
```
##########
File path: src/relay/op/contrib/ethosu/identity.cc
##########
@@ -0,0 +1,128 @@
+/*
+ * 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/op/contrib/ethosu/identity.cc
+ * \brief Property def of the Arm Ethos-U identity op.
+ */
+#include <tvm/relay/op.h>
+
+#include "common.h"
+
+namespace tvm {
+namespace relay {
+namespace op {
+namespace contrib {
+namespace ethosu {
+
+/*! \brief Attributes used by the Ethos(TM)-U NPU identity operator */
+struct EthosuIdentityAttrs : public tvm::AttrsNode<EthosuIdentityAttrs> {
+ double ifm_scale;
+ int ifm_zero_point;
+ double ofm_scale;
+ int ofm_zero_point;
+ String activation;
+
+ TVM_DECLARE_ATTRS(EthosuIdentityAttrs, "relay.attrs.EthosuIdentityAttrs") {
+ TVM_ATTR_FIELD(ifm_scale).describe("The quantization scale for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ifm_zero_point)
+ .describe("The quantization zero point for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_scale).describe("The quantization scale for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_zero_point)
+ .describe("The quantization zero point for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(activation)
+ .describe(
+ "The activation function to use. "
+ "'NONE' - no activation function. "
+ "'TANH' - tanh activation function. "
+ "'SIGMOID' - sigmoid activation function. "
+ "'LUT' - use a look-up table to perform the activation function.")
+ .set_default("NONE");
+ }
+};
+
+TVM_REGISTER_NODE_TYPE(EthosuIdentityAttrs);
+
+bool EthosuIdentityRel(const Array<Type>& types, int num_inputs, const Attrs& attrs,
+ const TypeReporter& reporter) {
+ int ifm_index = 0;
+ int result_index = 2;
+ ICHECK_EQ(types.size(), result_index + 1);
+
+ const auto* ifm = types[ifm_index].as<TensorTypeNode>();
+ if (ifm == nullptr) return false;
+
+ const auto* param = attrs.as<EthosuIdentityAttrs>();
+
+ ICHECK(param != nullptr) << "EthosuIdentityAttrs cannot be nullptr.";
+
+ if (ifm->dtype != DataType::UInt(8) && ifm->dtype != DataType::Int(8)) {
+ reporter->GetDiagCtx().EmitFatal(
+ Diagnostic::Error(reporter->GetSpan())
+ << "Invalid operator: Expected type(uint8) or type(int8) for ifm but was " << ifm->dtype);
+ return false;
+ }
+
+ if (ifm->shape.size() > 4) {
+ reporter->GetDiagCtx().EmitFatal(
+ Diagnostic::Error(reporter->GetSpan())
+ << "Invalid operator: Input Feature Map should be at most 4 dimensional, but was "
+ << ifm->shape);
+ return false;
+ }
+
+ // Assign ofm type
+ auto ofm_shape = ifm->shape;
+ reporter->Assign(types[result_index], TensorType(ofm_shape, ifm->dtype));
+ return true;
+}
+
+Expr MakeEthosuIdentity(Expr ifm, Expr lut, double ifm_scale, int ifm_zero_point, double ofm_scale,
+ int ofm_zero_point, String activation) {
+ auto attrs = make_object<EthosuIdentityAttrs>();
+ attrs->ifm_scale = ifm_scale;
+ attrs->ifm_zero_point = ifm_zero_point;
+ attrs->ofm_scale = ofm_scale;
+ attrs->ofm_zero_point = ofm_zero_point;
+ attrs->activation = std::move(activation);
+ static const Op& op = Op::Get("contrib.ethosu.identity");
+ return Call(op, {ifm, lut}, Attrs(attrs), {});
+}
+
+TVM_REGISTER_GLOBAL("relay.op._make.ethosu_identity").set_body_typed(MakeEthosuIdentity);
+
+RELAY_REGISTER_OP("contrib.ethosu.identity")
+ .describe(R"code(Identity operator for Ethos-U NPUs.
Review comment:
```suggestion
.describe(R"code(Arm(R) Ethos(TM)-U NPU identity operator.
```
##########
File path: src/relay/op/contrib/ethosu/identity.cc
##########
@@ -0,0 +1,128 @@
+/*
+ * 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/op/contrib/ethosu/identity.cc
+ * \brief Property def of the Arm Ethos-U identity op.
Review comment:
```suggestion
* \brief Identity operator definition for the Arm(R) Ethos(TM)-U NPU.
```
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] NicolaLancellotti commented on a change in pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
Posted by GitBox <gi...@apache.org>.
NicolaLancellotti commented on a change in pull request #9457:
URL: https://github.com/apache/tvm/pull/9457#discussion_r743598504
##########
File path: python/tvm/relay/backend/contrib/ethosu/te/dma.py
##########
@@ -67,63 +67,72 @@ def _pad(*indices):
return _pad
-def read_compute(tensor: te.Tensor, layout: str, zero_point: int, scale: float) -> te.Tensor:
+def read_compute(tensor: te.Tensor, zero_point: int, scale: float, layout: str = None) -> te.Tensor:
Review comment:
```suggestion
def read_compute(tensor: te.Tensor, zero_point: int, scale: float, layout: Optional[str] = None) -> te.Tensor:
```
##########
File path: src/relay/op/contrib/ethosu/identity.cc
##########
@@ -0,0 +1,128 @@
+/*
+ * 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/op/contrib/ethosu/identity.cc
+ * \brief Property def of the Arm Ethos-U identity op.
Review comment:
```suggestion
* \brief Identity operator definition for the Arm(R) Ethos(TM)-U NPU.
```
##########
File path: python/tvm/relay/backend/contrib/ethosu/op/identity.py
##########
@@ -0,0 +1,98 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+# pylint: disable=unused-argument
+"""Relay identity operator"""
Review comment:
```suggestion
"""Relay operator for identity for Arm(R) Ethos(TM)-U NPU"""
```
##########
File path: python/tvm/relay/backend/contrib/ethosu/legalize.py
##########
@@ -123,6 +123,108 @@ def __call__(self, *args, **kwargs):
pass
+class StridedSliceRewriter(DFPatternCallback):
+ """This pass brings the strided slice out of the partitioned function"""
+
+ def __init__(self):
+ super().__init__(require_type=True, rewrite_once=True)
+ self.pattern = (wildcard().has_attr({"Composite": "ethosu.strided_slice"}))(wildcard())
+
+ def callback(
+ self, pre: tvm.relay.Expr, post: tvm.relay.Expr, node_map: tvm.ir.container.Map
+ ) -> tvm.relay.Expr:
+ input = post.args[0]
+ attrs = post.op.body.attrs
+ begin = attrs.begin
+ end = attrs.end
+ strides = attrs.strides
+ axes = attrs.axes
+ slice_mode = attrs.slice_mode
+ strided_slice = relay.op.strided_slice(
+ input, begin, end, strides=strides, axes=axes, slice_mode=slice_mode
+ )
+ return strided_slice
+
+
+@ir.transform.module_pass(opt_level=1)
+class LegalizeStridedSlice:
+ """This is the pass that wraps StridedSliceRewriter"""
+
+ def transform_module(
+ self, mod: tvm.ir.IRModule, ctx: tvm.ir.transform.PassContext
+ ) -> tvm.ir.IRModule:
+ for global_var, func in mod.functions.items():
+ func = rewrite(StridedSliceRewriter(), func)
+ mod.update_func(global_var, func)
+ return mod
+
+ def __call__(self, *args, **kwargs):
+ pass
+
+
+class ReshapeRewriter(DFPatternCallback):
+ """This pass brings the reshape out of the partitioned function"""
+
+ def __init__(self):
+ super().__init__(require_type=True, rewrite_once=True)
+ self.pattern = (wildcard().has_attr({"Composite": "ethosu.reshape"}))(wildcard())
+
+ def callback(
+ self, pre: tvm.relay.Expr, post: tvm.relay.Expr, node_map: tvm.ir.container.Map
+ ) -> tvm.relay.Expr:
+ reshape_input = post.args[0]
+ new_shape = post.op.body.attrs.newshape
+ reshape = relay.op.reshape(reshape_input, newshape=new_shape)
+ return reshape
+
+
+@ir.transform.module_pass(opt_level=1)
+class LegalizeReshape:
+ """This is the pass that wraps ReshapeRewriter"""
+
+ def transform_module(
+ self, mod: tvm.ir.IRModule, ctx: tvm.ir.transform.PassContext
+ ) -> tvm.ir.IRModule:
+ for global_var, func in mod.functions.items():
+ func = rewrite(ReshapeRewriter(), func)
+ mod.update_func(global_var, func)
+ return mod
+
+ def __call__(self, *args, **kwargs):
+ pass
+
+
+class NoOpRewriter(DFPatternCallback):
+ """This pass adds and idenity operator to reshape and strided slice to avoid a no op without a consumer"""
Review comment:
```suggestion
"""This pass adds an identity operator to reshape and strided slice to avoid a no op without a consumer"""
```
##########
File path: src/relay/op/contrib/ethosu/identity.cc
##########
@@ -0,0 +1,128 @@
+/*
+ * 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/op/contrib/ethosu/identity.cc
+ * \brief Property def of the Arm Ethos-U identity op.
+ */
+#include <tvm/relay/op.h>
+
+#include "common.h"
+
+namespace tvm {
+namespace relay {
+namespace op {
+namespace contrib {
+namespace ethosu {
+
+/*! \brief Attributes used by the Ethos(TM)-U NPU identity operator */
+struct EthosuIdentityAttrs : public tvm::AttrsNode<EthosuIdentityAttrs> {
+ double ifm_scale;
+ int ifm_zero_point;
+ double ofm_scale;
+ int ofm_zero_point;
+ String activation;
+
+ TVM_DECLARE_ATTRS(EthosuIdentityAttrs, "relay.attrs.EthosuIdentityAttrs") {
+ TVM_ATTR_FIELD(ifm_scale).describe("The quantization scale for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ifm_zero_point)
+ .describe("The quantization zero point for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_scale).describe("The quantization scale for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_zero_point)
+ .describe("The quantization zero point for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(activation)
+ .describe(
+ "The activation function to use. "
+ "'NONE' - no activation function. "
+ "'TANH' - tanh activation function. "
+ "'SIGMOID' - sigmoid activation function. "
+ "'LUT' - use a look-up table to perform the activation function.")
+ .set_default("NONE");
+ }
+};
+
+TVM_REGISTER_NODE_TYPE(EthosuIdentityAttrs);
+
+bool EthosuIdentityRel(const Array<Type>& types, int num_inputs, const Attrs& attrs,
+ const TypeReporter& reporter) {
+ int ifm_index = 0;
+ int result_index = 2;
+ ICHECK_EQ(types.size(), result_index + 1);
+
+ const auto* ifm = types[ifm_index].as<TensorTypeNode>();
+ if (ifm == nullptr) return false;
+
+ const auto* param = attrs.as<EthosuIdentityAttrs>();
+
+ ICHECK(param != nullptr) << "EthosuIdentityAttrs cannot be nullptr.";
+
+ if (ifm->dtype != DataType::UInt(8) && ifm->dtype != DataType::Int(8)) {
+ reporter->GetDiagCtx().EmitFatal(
+ Diagnostic::Error(reporter->GetSpan())
+ << "Invalid operator: Expected type(uint8) or type(int8) for ifm but was " << ifm->dtype);
+ return false;
+ }
+
+ if (ifm->shape.size() > 4) {
+ reporter->GetDiagCtx().EmitFatal(
+ Diagnostic::Error(reporter->GetSpan())
+ << "Invalid operator: Input Feature Map should be at most 4 dimensional, but was "
+ << ifm->shape);
+ return false;
+ }
+
+ // Assign ofm type
+ auto ofm_shape = ifm->shape;
+ reporter->Assign(types[result_index], TensorType(ofm_shape, ifm->dtype));
+ return true;
+}
+
+Expr MakeEthosuIdentity(Expr ifm, Expr lut, double ifm_scale, int ifm_zero_point, double ofm_scale,
+ int ofm_zero_point, String activation) {
+ auto attrs = make_object<EthosuIdentityAttrs>();
+ attrs->ifm_scale = ifm_scale;
+ attrs->ifm_zero_point = ifm_zero_point;
+ attrs->ofm_scale = ofm_scale;
+ attrs->ofm_zero_point = ofm_zero_point;
+ attrs->activation = std::move(activation);
+ static const Op& op = Op::Get("contrib.ethosu.identity");
+ return Call(op, {ifm, lut}, Attrs(attrs), {});
+}
+
+TVM_REGISTER_GLOBAL("relay.op._make.ethosu_identity").set_body_typed(MakeEthosuIdentity);
+
+RELAY_REGISTER_OP("contrib.ethosu.identity")
+ .describe(R"code(Identity operator for Ethos-U NPUs.
Review comment:
```suggestion
.describe(R"code(Arm(R) Ethos(TM)-U NPU identity operator.
```
##########
File path: python/tvm/relay/backend/contrib/ethosu/te/dma.py
##########
@@ -67,63 +67,72 @@ def _pad(*indices):
return _pad
-def read_compute(tensor: te.Tensor, layout: str, zero_point: int, scale: float) -> te.Tensor:
+def read_compute(tensor: te.Tensor, zero_point: int, scale: float, layout: str = None) -> te.Tensor:
"""A tensor expression which represents a read.
Parameters
----------
tensor : te.Tensor
The tensor to read.
- layout : str
- The layout of the tensor, either NHWC or NHCWB16.
zero_point : int
The zero point of the tensor.
scale : float
The scale of the tensor.
+ layout : Optional[str]
+ The layout of the tensor, either NHWC or NHCWB16.
Returns
-------
te.Tensor
The tensor having been read.
"""
- assert layout in {"NHWC", "NHCWB16"}
read_attrs = {
"op": "ethosu_read",
- "layout": layout,
"zero_point": zero_point,
"scale": scale,
}
+
+ if layout:
+ assert layout in {"NHWC", "NHCWB16"}
+ read_attrs["layout"] = layout
+
return te.compute(tensor.shape, lambda *i: tensor(*i), name="ethosu_read", attrs=read_attrs)
-def write_compute(tensor: te.Tensor, layout: str, zero_point: int, scale: float) -> te.Tensor:
+def write_compute(
+ tensor: te.Tensor, zero_point: int, scale: float, layout: str = None
Review comment:
```suggestion
tensor: te.Tensor, zero_point: int, scale: float, layout: Optional[str] = None
```
##########
File path: python/tvm/relay/backend/contrib/ethosu/op/identity.py
##########
@@ -0,0 +1,98 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+# pylint: disable=unused-argument
+"""Relay identity operator"""
Review comment:
```suggestion
"""Relay operator for identity for Arm(R) Ethos(TM)-U NPU"""
```
##########
File path: src/relay/op/contrib/ethosu/identity.cc
##########
@@ -0,0 +1,128 @@
+/*
+ * 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/op/contrib/ethosu/identity.cc
+ * \brief Property def of the Arm Ethos-U identity op.
+ */
+#include <tvm/relay/op.h>
+
+#include "common.h"
+
+namespace tvm {
+namespace relay {
+namespace op {
+namespace contrib {
+namespace ethosu {
+
+/*! \brief Attributes used by the Ethos(TM)-U NPU identity operator */
+struct EthosuIdentityAttrs : public tvm::AttrsNode<EthosuIdentityAttrs> {
+ double ifm_scale;
+ int ifm_zero_point;
+ double ofm_scale;
+ int ofm_zero_point;
+ String activation;
+
+ TVM_DECLARE_ATTRS(EthosuIdentityAttrs, "relay.attrs.EthosuIdentityAttrs") {
+ TVM_ATTR_FIELD(ifm_scale).describe("The quantization scale for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ifm_zero_point)
+ .describe("The quantization zero point for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_scale).describe("The quantization scale for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_zero_point)
+ .describe("The quantization zero point for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(activation)
+ .describe(
+ "The activation function to use. "
+ "'NONE' - no activation function. "
+ "'TANH' - tanh activation function. "
+ "'SIGMOID' - sigmoid activation function. "
+ "'LUT' - use a look-up table to perform the activation function.")
+ .set_default("NONE");
+ }
+};
+
+TVM_REGISTER_NODE_TYPE(EthosuIdentityAttrs);
+
+bool EthosuIdentityRel(const Array<Type>& types, int num_inputs, const Attrs& attrs,
+ const TypeReporter& reporter) {
+ int ifm_index = 0;
+ int result_index = 2;
+ ICHECK_EQ(types.size(), result_index + 1);
+
+ const auto* ifm = types[ifm_index].as<TensorTypeNode>();
+ if (ifm == nullptr) return false;
+
+ const auto* param = attrs.as<EthosuIdentityAttrs>();
+
+ ICHECK(param != nullptr) << "EthosuIdentityAttrs cannot be nullptr.";
+
+ if (ifm->dtype != DataType::UInt(8) && ifm->dtype != DataType::Int(8)) {
+ reporter->GetDiagCtx().EmitFatal(
+ Diagnostic::Error(reporter->GetSpan())
+ << "Invalid operator: Expected type(uint8) or type(int8) for ifm but was " << ifm->dtype);
+ return false;
+ }
+
+ if (ifm->shape.size() > 4) {
+ reporter->GetDiagCtx().EmitFatal(
+ Diagnostic::Error(reporter->GetSpan())
+ << "Invalid operator: Input Feature Map should be at most 4 dimensional, but was "
+ << ifm->shape);
+ return false;
+ }
+
+ // Assign ofm type
+ auto ofm_shape = ifm->shape;
+ reporter->Assign(types[result_index], TensorType(ofm_shape, ifm->dtype));
+ return true;
+}
+
+Expr MakeEthosuIdentity(Expr ifm, Expr lut, double ifm_scale, int ifm_zero_point, double ofm_scale,
+ int ofm_zero_point, String activation) {
+ auto attrs = make_object<EthosuIdentityAttrs>();
+ attrs->ifm_scale = ifm_scale;
+ attrs->ifm_zero_point = ifm_zero_point;
+ attrs->ofm_scale = ofm_scale;
+ attrs->ofm_zero_point = ofm_zero_point;
+ attrs->activation = std::move(activation);
+ static const Op& op = Op::Get("contrib.ethosu.identity");
+ return Call(op, {ifm, lut}, Attrs(attrs), {});
+}
+
+TVM_REGISTER_GLOBAL("relay.op._make.ethosu_identity").set_body_typed(MakeEthosuIdentity);
+
+RELAY_REGISTER_OP("contrib.ethosu.identity")
+ .describe(R"code(Identity operator for Ethos-U NPUs.
Review comment:
```suggestion
.describe(R"code(Arm(R) Ethos(TM)-U NPU identity operator.
```
##########
File path: src/relay/op/contrib/ethosu/identity.cc
##########
@@ -0,0 +1,128 @@
+/*
+ * 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/op/contrib/ethosu/identity.cc
+ * \brief Property def of the Arm Ethos-U identity op.
Review comment:
```suggestion
* \brief Identity operator definition for the Arm(R) Ethos(TM)-U NPU.
```
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ekalda commented on a change in pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
Posted by GitBox <gi...@apache.org>.
ekalda commented on a change in pull request #9457:
URL: https://github.com/apache/tvm/pull/9457#discussion_r750115698
##########
File path: python/tvm/relay/backend/contrib/ethosu/te/dma.py
##########
@@ -67,63 +67,72 @@ def _pad(*indices):
return _pad
-def read_compute(tensor: te.Tensor, layout: str, zero_point: int, scale: float) -> te.Tensor:
+def read_compute(tensor: te.Tensor, zero_point: int, scale: float, layout: str = None) -> te.Tensor:
Review comment:
Done
##########
File path: python/tvm/relay/backend/contrib/ethosu/te/dma.py
##########
@@ -67,63 +67,72 @@ def _pad(*indices):
return _pad
-def read_compute(tensor: te.Tensor, layout: str, zero_point: int, scale: float) -> te.Tensor:
+def read_compute(tensor: te.Tensor, zero_point: int, scale: float, layout: str = None) -> te.Tensor:
"""A tensor expression which represents a read.
Parameters
----------
tensor : te.Tensor
The tensor to read.
- layout : str
- The layout of the tensor, either NHWC or NHCWB16.
zero_point : int
The zero point of the tensor.
scale : float
The scale of the tensor.
+ layout : Optional[str]
+ The layout of the tensor, either NHWC or NHCWB16.
Returns
-------
te.Tensor
The tensor having been read.
"""
- assert layout in {"NHWC", "NHCWB16"}
read_attrs = {
"op": "ethosu_read",
- "layout": layout,
"zero_point": zero_point,
"scale": scale,
}
+
+ if layout:
+ assert layout in {"NHWC", "NHCWB16"}
+ read_attrs["layout"] = layout
+
return te.compute(tensor.shape, lambda *i: tensor(*i), name="ethosu_read", attrs=read_attrs)
-def write_compute(tensor: te.Tensor, layout: str, zero_point: int, scale: float) -> te.Tensor:
+def write_compute(
+ tensor: te.Tensor, zero_point: int, scale: float, layout: str = None
Review comment:
Done
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ekalda commented on a change in pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
Posted by GitBox <gi...@apache.org>.
ekalda commented on a change in pull request #9457:
URL: https://github.com/apache/tvm/pull/9457#discussion_r750119798
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir/identity.py
##########
@@ -0,0 +1,87 @@
+# 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
+"""Extract information from the identity operator in TIR."""
+from typing import Dict, Tuple
+import tvm
+from .dma import get_read_params, get_write_params
+from .spec import SerialKernel, SerialActivation, SerialPooling, SerialPadding
+from .utils import get_op_attrs
+
+
+def get_identity_params(
+ stmt: tvm.tir.AttrStmt,
+ producers: Dict[tvm.tir.Var, tvm.tir.AttrStmt],
+ consumers: Dict[tvm.tir.Var, tvm.tir.AttrStmt],
+) -> Tuple[SerialPooling, tvm.tir.Var, tvm.tir.Var]:
+ """Get the parameters necessary to construct a call_extern for a pooling.
+
+ Parameters
+ ----------
+ stmt : tvm.tir.AttrStmt
+ The outermost attribute statement of a convolution loop nest.
Review comment:
Done
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ekalda commented on a change in pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
Posted by GitBox <gi...@apache.org>.
ekalda commented on a change in pull request #9457:
URL: https://github.com/apache/tvm/pull/9457#discussion_r750118808
##########
File path: src/relay/op/contrib/ethosu/identity.cc
##########
@@ -0,0 +1,128 @@
+/*
+ * 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/op/contrib/ethosu/identity.cc
+ * \brief Property def of the Arm Ethos-U identity op.
+ */
+#include <tvm/relay/op.h>
+
+#include "common.h"
+
+namespace tvm {
+namespace relay {
+namespace op {
+namespace contrib {
+namespace ethosu {
+
+/*! \brief Attributes used by the Ethos(TM)-U NPU identity operator */
+struct EthosuIdentityAttrs : public tvm::AttrsNode<EthosuIdentityAttrs> {
+ double ifm_scale;
+ int ifm_zero_point;
+ double ofm_scale;
+ int ofm_zero_point;
+ String activation;
+
+ TVM_DECLARE_ATTRS(EthosuIdentityAttrs, "relay.attrs.EthosuIdentityAttrs") {
+ TVM_ATTR_FIELD(ifm_scale).describe("The quantization scale for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ifm_zero_point)
+ .describe("The quantization zero point for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_scale).describe("The quantization scale for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_zero_point)
+ .describe("The quantization zero point for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(activation)
+ .describe(
+ "The activation function to use. "
+ "'NONE' - no activation function. "
+ "'TANH' - tanh activation function. "
+ "'SIGMOID' - sigmoid activation function. "
+ "'LUT' - use a look-up table to perform the activation function.")
+ .set_default("NONE");
+ }
+};
+
+TVM_REGISTER_NODE_TYPE(EthosuIdentityAttrs);
+
+bool EthosuIdentityRel(const Array<Type>& types, int num_inputs, const Attrs& attrs,
+ const TypeReporter& reporter) {
+ int ifm_index = 0;
+ int result_index = 2;
+ ICHECK_EQ(types.size(), result_index + 1);
+
+ const auto* ifm = types[ifm_index].as<TensorTypeNode>();
+ if (ifm == nullptr) return false;
+
+ const auto* param = attrs.as<EthosuIdentityAttrs>();
+
+ ICHECK(param != nullptr) << "EthosuIdentityAttrs cannot be nullptr.";
+
+ if (ifm->dtype != DataType::UInt(8) && ifm->dtype != DataType::Int(8)) {
+ reporter->GetDiagCtx().EmitFatal(
+ Diagnostic::Error(reporter->GetSpan())
+ << "Invalid operator: Expected type(uint8) or type(int8) for ifm but was " << ifm->dtype);
+ return false;
+ }
+
+ if (ifm->shape.size() > 4) {
+ reporter->GetDiagCtx().EmitFatal(
+ Diagnostic::Error(reporter->GetSpan())
+ << "Invalid operator: Input Feature Map should be at most 4 dimensional, but was "
+ << ifm->shape);
+ return false;
+ }
+
+ // Assign ofm type
+ auto ofm_shape = ifm->shape;
+ reporter->Assign(types[result_index], TensorType(ofm_shape, ifm->dtype));
+ return true;
+}
+
+Expr MakeEthosuIdentity(Expr ifm, Expr lut, double ifm_scale, int ifm_zero_point, double ofm_scale,
+ int ofm_zero_point, String activation) {
+ auto attrs = make_object<EthosuIdentityAttrs>();
+ attrs->ifm_scale = ifm_scale;
+ attrs->ifm_zero_point = ifm_zero_point;
+ attrs->ofm_scale = ofm_scale;
+ attrs->ofm_zero_point = ofm_zero_point;
+ attrs->activation = std::move(activation);
+ static const Op& op = Op::Get("contrib.ethosu.identity");
+ return Call(op, {ifm, lut}, Attrs(attrs), {});
+}
+
+TVM_REGISTER_GLOBAL("relay.op._make.ethosu_identity").set_body_typed(MakeEthosuIdentity);
+
+RELAY_REGISTER_OP("contrib.ethosu.identity")
+ .describe(R"code(Identity operator for Ethos-U NPUs.
+
+This Relay operator performs the identity operation on Ethos(TM)-U NPU with a capability
+to requantize the data. It accepts input with any shape that is less or equal to 4.
Review comment:
Done
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] mbaret merged pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
Posted by GitBox <gi...@apache.org>.
mbaret merged pull request #9457:
URL: https://github.com/apache/tvm/pull/9457
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] NicolaLancellotti commented on a change in pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
Posted by GitBox <gi...@apache.org>.
NicolaLancellotti commented on a change in pull request #9457:
URL: https://github.com/apache/tvm/pull/9457#discussion_r743598504
##########
File path: python/tvm/relay/backend/contrib/ethosu/te/dma.py
##########
@@ -67,63 +67,72 @@ def _pad(*indices):
return _pad
-def read_compute(tensor: te.Tensor, layout: str, zero_point: int, scale: float) -> te.Tensor:
+def read_compute(tensor: te.Tensor, zero_point: int, scale: float, layout: str = None) -> te.Tensor:
Review comment:
```suggestion
def read_compute(tensor: te.Tensor, zero_point: int, scale: float, layout: Optional[str] = None) -> te.Tensor:
```
##########
File path: src/relay/op/contrib/ethosu/identity.cc
##########
@@ -0,0 +1,128 @@
+/*
+ * 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/op/contrib/ethosu/identity.cc
+ * \brief Property def of the Arm Ethos-U identity op.
Review comment:
```suggestion
* \brief Identity operator definition for the Arm(R) Ethos(TM)-U NPU.
```
##########
File path: python/tvm/relay/backend/contrib/ethosu/op/identity.py
##########
@@ -0,0 +1,98 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+# pylint: disable=unused-argument
+"""Relay identity operator"""
Review comment:
```suggestion
"""Relay operator for identity for Arm(R) Ethos(TM)-U NPU"""
```
##########
File path: python/tvm/relay/backend/contrib/ethosu/legalize.py
##########
@@ -123,6 +123,108 @@ def __call__(self, *args, **kwargs):
pass
+class StridedSliceRewriter(DFPatternCallback):
+ """This pass brings the strided slice out of the partitioned function"""
+
+ def __init__(self):
+ super().__init__(require_type=True, rewrite_once=True)
+ self.pattern = (wildcard().has_attr({"Composite": "ethosu.strided_slice"}))(wildcard())
+
+ def callback(
+ self, pre: tvm.relay.Expr, post: tvm.relay.Expr, node_map: tvm.ir.container.Map
+ ) -> tvm.relay.Expr:
+ input = post.args[0]
+ attrs = post.op.body.attrs
+ begin = attrs.begin
+ end = attrs.end
+ strides = attrs.strides
+ axes = attrs.axes
+ slice_mode = attrs.slice_mode
+ strided_slice = relay.op.strided_slice(
+ input, begin, end, strides=strides, axes=axes, slice_mode=slice_mode
+ )
+ return strided_slice
+
+
+@ir.transform.module_pass(opt_level=1)
+class LegalizeStridedSlice:
+ """This is the pass that wraps StridedSliceRewriter"""
+
+ def transform_module(
+ self, mod: tvm.ir.IRModule, ctx: tvm.ir.transform.PassContext
+ ) -> tvm.ir.IRModule:
+ for global_var, func in mod.functions.items():
+ func = rewrite(StridedSliceRewriter(), func)
+ mod.update_func(global_var, func)
+ return mod
+
+ def __call__(self, *args, **kwargs):
+ pass
+
+
+class ReshapeRewriter(DFPatternCallback):
+ """This pass brings the reshape out of the partitioned function"""
+
+ def __init__(self):
+ super().__init__(require_type=True, rewrite_once=True)
+ self.pattern = (wildcard().has_attr({"Composite": "ethosu.reshape"}))(wildcard())
+
+ def callback(
+ self, pre: tvm.relay.Expr, post: tvm.relay.Expr, node_map: tvm.ir.container.Map
+ ) -> tvm.relay.Expr:
+ reshape_input = post.args[0]
+ new_shape = post.op.body.attrs.newshape
+ reshape = relay.op.reshape(reshape_input, newshape=new_shape)
+ return reshape
+
+
+@ir.transform.module_pass(opt_level=1)
+class LegalizeReshape:
+ """This is the pass that wraps ReshapeRewriter"""
+
+ def transform_module(
+ self, mod: tvm.ir.IRModule, ctx: tvm.ir.transform.PassContext
+ ) -> tvm.ir.IRModule:
+ for global_var, func in mod.functions.items():
+ func = rewrite(ReshapeRewriter(), func)
+ mod.update_func(global_var, func)
+ return mod
+
+ def __call__(self, *args, **kwargs):
+ pass
+
+
+class NoOpRewriter(DFPatternCallback):
+ """This pass adds and idenity operator to reshape and strided slice to avoid a no op without a consumer"""
Review comment:
```suggestion
"""This pass adds an identity operator to reshape and strided slice to avoid a no op without a consumer"""
```
##########
File path: src/relay/op/contrib/ethosu/identity.cc
##########
@@ -0,0 +1,128 @@
+/*
+ * 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/op/contrib/ethosu/identity.cc
+ * \brief Property def of the Arm Ethos-U identity op.
+ */
+#include <tvm/relay/op.h>
+
+#include "common.h"
+
+namespace tvm {
+namespace relay {
+namespace op {
+namespace contrib {
+namespace ethosu {
+
+/*! \brief Attributes used by the Ethos(TM)-U NPU identity operator */
+struct EthosuIdentityAttrs : public tvm::AttrsNode<EthosuIdentityAttrs> {
+ double ifm_scale;
+ int ifm_zero_point;
+ double ofm_scale;
+ int ofm_zero_point;
+ String activation;
+
+ TVM_DECLARE_ATTRS(EthosuIdentityAttrs, "relay.attrs.EthosuIdentityAttrs") {
+ TVM_ATTR_FIELD(ifm_scale).describe("The quantization scale for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ifm_zero_point)
+ .describe("The quantization zero point for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_scale).describe("The quantization scale for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_zero_point)
+ .describe("The quantization zero point for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(activation)
+ .describe(
+ "The activation function to use. "
+ "'NONE' - no activation function. "
+ "'TANH' - tanh activation function. "
+ "'SIGMOID' - sigmoid activation function. "
+ "'LUT' - use a look-up table to perform the activation function.")
+ .set_default("NONE");
+ }
+};
+
+TVM_REGISTER_NODE_TYPE(EthosuIdentityAttrs);
+
+bool EthosuIdentityRel(const Array<Type>& types, int num_inputs, const Attrs& attrs,
+ const TypeReporter& reporter) {
+ int ifm_index = 0;
+ int result_index = 2;
+ ICHECK_EQ(types.size(), result_index + 1);
+
+ const auto* ifm = types[ifm_index].as<TensorTypeNode>();
+ if (ifm == nullptr) return false;
+
+ const auto* param = attrs.as<EthosuIdentityAttrs>();
+
+ ICHECK(param != nullptr) << "EthosuIdentityAttrs cannot be nullptr.";
+
+ if (ifm->dtype != DataType::UInt(8) && ifm->dtype != DataType::Int(8)) {
+ reporter->GetDiagCtx().EmitFatal(
+ Diagnostic::Error(reporter->GetSpan())
+ << "Invalid operator: Expected type(uint8) or type(int8) for ifm but was " << ifm->dtype);
+ return false;
+ }
+
+ if (ifm->shape.size() > 4) {
+ reporter->GetDiagCtx().EmitFatal(
+ Diagnostic::Error(reporter->GetSpan())
+ << "Invalid operator: Input Feature Map should be at most 4 dimensional, but was "
+ << ifm->shape);
+ return false;
+ }
+
+ // Assign ofm type
+ auto ofm_shape = ifm->shape;
+ reporter->Assign(types[result_index], TensorType(ofm_shape, ifm->dtype));
+ return true;
+}
+
+Expr MakeEthosuIdentity(Expr ifm, Expr lut, double ifm_scale, int ifm_zero_point, double ofm_scale,
+ int ofm_zero_point, String activation) {
+ auto attrs = make_object<EthosuIdentityAttrs>();
+ attrs->ifm_scale = ifm_scale;
+ attrs->ifm_zero_point = ifm_zero_point;
+ attrs->ofm_scale = ofm_scale;
+ attrs->ofm_zero_point = ofm_zero_point;
+ attrs->activation = std::move(activation);
+ static const Op& op = Op::Get("contrib.ethosu.identity");
+ return Call(op, {ifm, lut}, Attrs(attrs), {});
+}
+
+TVM_REGISTER_GLOBAL("relay.op._make.ethosu_identity").set_body_typed(MakeEthosuIdentity);
+
+RELAY_REGISTER_OP("contrib.ethosu.identity")
+ .describe(R"code(Identity operator for Ethos-U NPUs.
Review comment:
```suggestion
.describe(R"code(Arm(R) Ethos(TM)-U NPU identity operator.
```
##########
File path: python/tvm/relay/backend/contrib/ethosu/te/dma.py
##########
@@ -67,63 +67,72 @@ def _pad(*indices):
return _pad
-def read_compute(tensor: te.Tensor, layout: str, zero_point: int, scale: float) -> te.Tensor:
+def read_compute(tensor: te.Tensor, zero_point: int, scale: float, layout: str = None) -> te.Tensor:
"""A tensor expression which represents a read.
Parameters
----------
tensor : te.Tensor
The tensor to read.
- layout : str
- The layout of the tensor, either NHWC or NHCWB16.
zero_point : int
The zero point of the tensor.
scale : float
The scale of the tensor.
+ layout : Optional[str]
+ The layout of the tensor, either NHWC or NHCWB16.
Returns
-------
te.Tensor
The tensor having been read.
"""
- assert layout in {"NHWC", "NHCWB16"}
read_attrs = {
"op": "ethosu_read",
- "layout": layout,
"zero_point": zero_point,
"scale": scale,
}
+
+ if layout:
+ assert layout in {"NHWC", "NHCWB16"}
+ read_attrs["layout"] = layout
+
return te.compute(tensor.shape, lambda *i: tensor(*i), name="ethosu_read", attrs=read_attrs)
-def write_compute(tensor: te.Tensor, layout: str, zero_point: int, scale: float) -> te.Tensor:
+def write_compute(
+ tensor: te.Tensor, zero_point: int, scale: float, layout: str = None
Review comment:
```suggestion
tensor: te.Tensor, zero_point: int, scale: float, layout: Optional[str] = None
```
##########
File path: python/tvm/relay/backend/contrib/ethosu/op/identity.py
##########
@@ -0,0 +1,98 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+# pylint: disable=unused-argument
+"""Relay identity operator"""
Review comment:
```suggestion
"""Relay operator for identity for Arm(R) Ethos(TM)-U NPU"""
```
##########
File path: src/relay/op/contrib/ethosu/identity.cc
##########
@@ -0,0 +1,128 @@
+/*
+ * 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/op/contrib/ethosu/identity.cc
+ * \brief Property def of the Arm Ethos-U identity op.
+ */
+#include <tvm/relay/op.h>
+
+#include "common.h"
+
+namespace tvm {
+namespace relay {
+namespace op {
+namespace contrib {
+namespace ethosu {
+
+/*! \brief Attributes used by the Ethos(TM)-U NPU identity operator */
+struct EthosuIdentityAttrs : public tvm::AttrsNode<EthosuIdentityAttrs> {
+ double ifm_scale;
+ int ifm_zero_point;
+ double ofm_scale;
+ int ofm_zero_point;
+ String activation;
+
+ TVM_DECLARE_ATTRS(EthosuIdentityAttrs, "relay.attrs.EthosuIdentityAttrs") {
+ TVM_ATTR_FIELD(ifm_scale).describe("The quantization scale for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ifm_zero_point)
+ .describe("The quantization zero point for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_scale).describe("The quantization scale for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_zero_point)
+ .describe("The quantization zero point for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(activation)
+ .describe(
+ "The activation function to use. "
+ "'NONE' - no activation function. "
+ "'TANH' - tanh activation function. "
+ "'SIGMOID' - sigmoid activation function. "
+ "'LUT' - use a look-up table to perform the activation function.")
+ .set_default("NONE");
+ }
+};
+
+TVM_REGISTER_NODE_TYPE(EthosuIdentityAttrs);
+
+bool EthosuIdentityRel(const Array<Type>& types, int num_inputs, const Attrs& attrs,
+ const TypeReporter& reporter) {
+ int ifm_index = 0;
+ int result_index = 2;
+ ICHECK_EQ(types.size(), result_index + 1);
+
+ const auto* ifm = types[ifm_index].as<TensorTypeNode>();
+ if (ifm == nullptr) return false;
+
+ const auto* param = attrs.as<EthosuIdentityAttrs>();
+
+ ICHECK(param != nullptr) << "EthosuIdentityAttrs cannot be nullptr.";
+
+ if (ifm->dtype != DataType::UInt(8) && ifm->dtype != DataType::Int(8)) {
+ reporter->GetDiagCtx().EmitFatal(
+ Diagnostic::Error(reporter->GetSpan())
+ << "Invalid operator: Expected type(uint8) or type(int8) for ifm but was " << ifm->dtype);
+ return false;
+ }
+
+ if (ifm->shape.size() > 4) {
+ reporter->GetDiagCtx().EmitFatal(
+ Diagnostic::Error(reporter->GetSpan())
+ << "Invalid operator: Input Feature Map should be at most 4 dimensional, but was "
+ << ifm->shape);
+ return false;
+ }
+
+ // Assign ofm type
+ auto ofm_shape = ifm->shape;
+ reporter->Assign(types[result_index], TensorType(ofm_shape, ifm->dtype));
+ return true;
+}
+
+Expr MakeEthosuIdentity(Expr ifm, Expr lut, double ifm_scale, int ifm_zero_point, double ofm_scale,
+ int ofm_zero_point, String activation) {
+ auto attrs = make_object<EthosuIdentityAttrs>();
+ attrs->ifm_scale = ifm_scale;
+ attrs->ifm_zero_point = ifm_zero_point;
+ attrs->ofm_scale = ofm_scale;
+ attrs->ofm_zero_point = ofm_zero_point;
+ attrs->activation = std::move(activation);
+ static const Op& op = Op::Get("contrib.ethosu.identity");
+ return Call(op, {ifm, lut}, Attrs(attrs), {});
+}
+
+TVM_REGISTER_GLOBAL("relay.op._make.ethosu_identity").set_body_typed(MakeEthosuIdentity);
+
+RELAY_REGISTER_OP("contrib.ethosu.identity")
+ .describe(R"code(Identity operator for Ethos-U NPUs.
Review comment:
```suggestion
.describe(R"code(Arm(R) Ethos(TM)-U NPU identity operator.
```
##########
File path: src/relay/op/contrib/ethosu/identity.cc
##########
@@ -0,0 +1,128 @@
+/*
+ * 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/op/contrib/ethosu/identity.cc
+ * \brief Property def of the Arm Ethos-U identity op.
Review comment:
```suggestion
* \brief Identity operator definition for the Arm(R) Ethos(TM)-U NPU.
```
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ekalda commented on a change in pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
Posted by GitBox <gi...@apache.org>.
ekalda commented on a change in pull request #9457:
URL: https://github.com/apache/tvm/pull/9457#discussion_r750120691
##########
File path: python/tvm/relay/backend/contrib/ethosu/legalize.py
##########
@@ -123,6 +123,109 @@ def __call__(self, *args, **kwargs):
pass
+class StridedSliceRewriter(DFPatternCallback):
+ """This pass brings the strided slice out of the partitioned function"""
+
+ def __init__(self):
+ super().__init__(require_type=True, rewrite_once=True)
+ self.pattern = (wildcard().has_attr({"Composite": "ethosu.strided_slice"}))(wildcard())
+
+ def callback(
+ self, pre: tvm.relay.Expr, post: tvm.relay.Expr, node_map: tvm.ir.container.Map
+ ) -> tvm.relay.Expr:
+ slice_input = post.args[0]
+ attrs = post.op.body.attrs
+ begin = attrs.begin
+ end = attrs.end
+ strides = attrs.strides
+ axes = attrs.axes
+ slice_mode = attrs.slice_mode
+ strided_slice = relay.op.strided_slice(
+ slice_input, begin, end, strides=strides, axes=axes, slice_mode=slice_mode
+ )
+ return strided_slice
+
+
+@ir.transform.module_pass(opt_level=1)
+class LegalizeStridedSlice:
+ """This is the pass that wraps StridedSliceRewriter"""
+
+ def transform_module(
+ self, mod: tvm.ir.IRModule, ctx: tvm.ir.transform.PassContext
+ ) -> tvm.ir.IRModule:
+ for global_var, func in mod.functions.items():
+ func = rewrite(StridedSliceRewriter(), func)
+ mod.update_func(global_var, func)
+ return mod
+
+ def __call__(self, *args, **kwargs):
+ pass
+
+
+class ReshapeRewriter(DFPatternCallback):
+ """This pass brings the reshape out of the partitioned function"""
+
+ def __init__(self):
+ super().__init__(require_type=True, rewrite_once=True)
+ self.pattern = (wildcard().has_attr({"Composite": "ethosu.reshape"}))(wildcard())
+
+ def callback(
+ self, pre: tvm.relay.Expr, post: tvm.relay.Expr, node_map: tvm.ir.container.Map
+ ) -> tvm.relay.Expr:
+ reshape_input = post.args[0]
+ new_shape = post.op.body.attrs.newshape
+ reshape = relay.op.reshape(reshape_input, newshape=new_shape)
+ return reshape
+
+
+@ir.transform.module_pass(opt_level=1)
+class LegalizeReshape:
+ """This is the pass that wraps ReshapeRewriter"""
+
+ def transform_module(
+ self, mod: tvm.ir.IRModule, ctx: tvm.ir.transform.PassContext
+ ) -> tvm.ir.IRModule:
+ for global_var, func in mod.functions.items():
+ func = rewrite(ReshapeRewriter(), func)
+ mod.update_func(global_var, func)
+ return mod
+
+ def __call__(self, *args, **kwargs):
+ pass
+
+
+class NoOpRewriter(DFPatternCallback):
+ """This pass adds and idenity operator to reshape and strided slice to avoid a no op
Review comment:
Done
##########
File path: python/tvm/relay/backend/contrib/ethosu/te/identity.py
##########
@@ -0,0 +1,73 @@
+# 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
+"""Tensor Expression for identity"""
+from tvm import te
+from .dma import read_compute, write_compute
+
+
+def identity_compute(
+ ifm: te.Tensor,
+ lut: te.Tensor,
+ ifm_scale: float,
+ ifm_zero_point: int,
+ ofm_scale: float,
+ ofm_zero_point: int,
+ activation: str,
+) -> te.Tensor:
+ """A compute operator for the NPU identity operator.
+
+ Parameters
+ ----------
+ ifm : te.Tensor
+ The Input Feature Map tensor (IFM).
+ lut : te.Tensor
+ The look-up table values to use if activation is "LUT", "TANH" or "SIGMOID".
+ ifm_scale : float
+ The quantization scale for the Input Feature Map tensor.
+ ifm_zero_point : int
+ The quantization zero point for the Input Feature Map tensor.
+ ofm_scale : float
+ The quantization scale for the Output Feature Map tensor.
+ ofm_zero_point : int
+ The quantization zero point for the Output Feature Map tensor.
+ activation : str
+ The activation function to use.
+ "NONE" - no activation function.
+ "TANH" - tanh activation function.
+ "SIGMOID" - sigmoid activation function.
+ "LUT" - use a look-up table to perform the activation function.
+
+ Returns
+ -------
+ te.Tensor
+ The OFM tensor.
Review comment:
Done
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ekalda commented on a change in pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
Posted by GitBox <gi...@apache.org>.
ekalda commented on a change in pull request #9457:
URL: https://github.com/apache/tvm/pull/9457#discussion_r750121349
##########
File path: tests/python/contrib/test_ethosu/test_type_inference.py
##########
@@ -14,6 +14,14 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
+from tvm import relay, TVMError
+from tvm import relay
+import tvm
+from tvm.relay.testing import run_opt_pass
+from .infra import make_ethosu_conv2d
+from .infra import make_ethosu_pooling
+from .infra import make_ethosu_identity
+
Review comment:
Oops... removed the double imports!
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ekalda commented on pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
Posted by GitBox <gi...@apache.org>.
ekalda commented on pull request #9457:
URL: https://github.com/apache/tvm/pull/9457#issuecomment-970128409
@mbaret @NicolaLancellotti @lhutton1 thanks for the reviews, much appreciated! :)
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] NicolaLancellotti commented on a change in pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
Posted by GitBox <gi...@apache.org>.
NicolaLancellotti commented on a change in pull request #9457:
URL: https://github.com/apache/tvm/pull/9457#discussion_r743598504
##########
File path: python/tvm/relay/backend/contrib/ethosu/te/dma.py
##########
@@ -67,63 +67,72 @@ def _pad(*indices):
return _pad
-def read_compute(tensor: te.Tensor, layout: str, zero_point: int, scale: float) -> te.Tensor:
+def read_compute(tensor: te.Tensor, zero_point: int, scale: float, layout: str = None) -> te.Tensor:
Review comment:
```suggestion
def read_compute(tensor: te.Tensor, zero_point: int, scale: float, layout: Optional[str] = None) -> te.Tensor:
```
##########
File path: src/relay/op/contrib/ethosu/identity.cc
##########
@@ -0,0 +1,128 @@
+/*
+ * 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/op/contrib/ethosu/identity.cc
+ * \brief Property def of the Arm Ethos-U identity op.
Review comment:
```suggestion
* \brief Identity operator definition for the Arm(R) Ethos(TM)-U NPU.
```
##########
File path: python/tvm/relay/backend/contrib/ethosu/op/identity.py
##########
@@ -0,0 +1,98 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+# pylint: disable=unused-argument
+"""Relay identity operator"""
Review comment:
```suggestion
"""Relay operator for identity for Arm(R) Ethos(TM)-U NPU"""
```
##########
File path: python/tvm/relay/backend/contrib/ethosu/legalize.py
##########
@@ -123,6 +123,108 @@ def __call__(self, *args, **kwargs):
pass
+class StridedSliceRewriter(DFPatternCallback):
+ """This pass brings the strided slice out of the partitioned function"""
+
+ def __init__(self):
+ super().__init__(require_type=True, rewrite_once=True)
+ self.pattern = (wildcard().has_attr({"Composite": "ethosu.strided_slice"}))(wildcard())
+
+ def callback(
+ self, pre: tvm.relay.Expr, post: tvm.relay.Expr, node_map: tvm.ir.container.Map
+ ) -> tvm.relay.Expr:
+ input = post.args[0]
+ attrs = post.op.body.attrs
+ begin = attrs.begin
+ end = attrs.end
+ strides = attrs.strides
+ axes = attrs.axes
+ slice_mode = attrs.slice_mode
+ strided_slice = relay.op.strided_slice(
+ input, begin, end, strides=strides, axes=axes, slice_mode=slice_mode
+ )
+ return strided_slice
+
+
+@ir.transform.module_pass(opt_level=1)
+class LegalizeStridedSlice:
+ """This is the pass that wraps StridedSliceRewriter"""
+
+ def transform_module(
+ self, mod: tvm.ir.IRModule, ctx: tvm.ir.transform.PassContext
+ ) -> tvm.ir.IRModule:
+ for global_var, func in mod.functions.items():
+ func = rewrite(StridedSliceRewriter(), func)
+ mod.update_func(global_var, func)
+ return mod
+
+ def __call__(self, *args, **kwargs):
+ pass
+
+
+class ReshapeRewriter(DFPatternCallback):
+ """This pass brings the reshape out of the partitioned function"""
+
+ def __init__(self):
+ super().__init__(require_type=True, rewrite_once=True)
+ self.pattern = (wildcard().has_attr({"Composite": "ethosu.reshape"}))(wildcard())
+
+ def callback(
+ self, pre: tvm.relay.Expr, post: tvm.relay.Expr, node_map: tvm.ir.container.Map
+ ) -> tvm.relay.Expr:
+ reshape_input = post.args[0]
+ new_shape = post.op.body.attrs.newshape
+ reshape = relay.op.reshape(reshape_input, newshape=new_shape)
+ return reshape
+
+
+@ir.transform.module_pass(opt_level=1)
+class LegalizeReshape:
+ """This is the pass that wraps ReshapeRewriter"""
+
+ def transform_module(
+ self, mod: tvm.ir.IRModule, ctx: tvm.ir.transform.PassContext
+ ) -> tvm.ir.IRModule:
+ for global_var, func in mod.functions.items():
+ func = rewrite(ReshapeRewriter(), func)
+ mod.update_func(global_var, func)
+ return mod
+
+ def __call__(self, *args, **kwargs):
+ pass
+
+
+class NoOpRewriter(DFPatternCallback):
+ """This pass adds and idenity operator to reshape and strided slice to avoid a no op without a consumer"""
Review comment:
```suggestion
"""This pass adds an identity operator to reshape and strided slice to avoid a no op without a consumer"""
```
##########
File path: src/relay/op/contrib/ethosu/identity.cc
##########
@@ -0,0 +1,128 @@
+/*
+ * 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/op/contrib/ethosu/identity.cc
+ * \brief Property def of the Arm Ethos-U identity op.
+ */
+#include <tvm/relay/op.h>
+
+#include "common.h"
+
+namespace tvm {
+namespace relay {
+namespace op {
+namespace contrib {
+namespace ethosu {
+
+/*! \brief Attributes used by the Ethos(TM)-U NPU identity operator */
+struct EthosuIdentityAttrs : public tvm::AttrsNode<EthosuIdentityAttrs> {
+ double ifm_scale;
+ int ifm_zero_point;
+ double ofm_scale;
+ int ofm_zero_point;
+ String activation;
+
+ TVM_DECLARE_ATTRS(EthosuIdentityAttrs, "relay.attrs.EthosuIdentityAttrs") {
+ TVM_ATTR_FIELD(ifm_scale).describe("The quantization scale for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ifm_zero_point)
+ .describe("The quantization zero point for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_scale).describe("The quantization scale for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_zero_point)
+ .describe("The quantization zero point for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(activation)
+ .describe(
+ "The activation function to use. "
+ "'NONE' - no activation function. "
+ "'TANH' - tanh activation function. "
+ "'SIGMOID' - sigmoid activation function. "
+ "'LUT' - use a look-up table to perform the activation function.")
+ .set_default("NONE");
+ }
+};
+
+TVM_REGISTER_NODE_TYPE(EthosuIdentityAttrs);
+
+bool EthosuIdentityRel(const Array<Type>& types, int num_inputs, const Attrs& attrs,
+ const TypeReporter& reporter) {
+ int ifm_index = 0;
+ int result_index = 2;
+ ICHECK_EQ(types.size(), result_index + 1);
+
+ const auto* ifm = types[ifm_index].as<TensorTypeNode>();
+ if (ifm == nullptr) return false;
+
+ const auto* param = attrs.as<EthosuIdentityAttrs>();
+
+ ICHECK(param != nullptr) << "EthosuIdentityAttrs cannot be nullptr.";
+
+ if (ifm->dtype != DataType::UInt(8) && ifm->dtype != DataType::Int(8)) {
+ reporter->GetDiagCtx().EmitFatal(
+ Diagnostic::Error(reporter->GetSpan())
+ << "Invalid operator: Expected type(uint8) or type(int8) for ifm but was " << ifm->dtype);
+ return false;
+ }
+
+ if (ifm->shape.size() > 4) {
+ reporter->GetDiagCtx().EmitFatal(
+ Diagnostic::Error(reporter->GetSpan())
+ << "Invalid operator: Input Feature Map should be at most 4 dimensional, but was "
+ << ifm->shape);
+ return false;
+ }
+
+ // Assign ofm type
+ auto ofm_shape = ifm->shape;
+ reporter->Assign(types[result_index], TensorType(ofm_shape, ifm->dtype));
+ return true;
+}
+
+Expr MakeEthosuIdentity(Expr ifm, Expr lut, double ifm_scale, int ifm_zero_point, double ofm_scale,
+ int ofm_zero_point, String activation) {
+ auto attrs = make_object<EthosuIdentityAttrs>();
+ attrs->ifm_scale = ifm_scale;
+ attrs->ifm_zero_point = ifm_zero_point;
+ attrs->ofm_scale = ofm_scale;
+ attrs->ofm_zero_point = ofm_zero_point;
+ attrs->activation = std::move(activation);
+ static const Op& op = Op::Get("contrib.ethosu.identity");
+ return Call(op, {ifm, lut}, Attrs(attrs), {});
+}
+
+TVM_REGISTER_GLOBAL("relay.op._make.ethosu_identity").set_body_typed(MakeEthosuIdentity);
+
+RELAY_REGISTER_OP("contrib.ethosu.identity")
+ .describe(R"code(Identity operator for Ethos-U NPUs.
Review comment:
```suggestion
.describe(R"code(Arm(R) Ethos(TM)-U NPU identity operator.
```
##########
File path: python/tvm/relay/backend/contrib/ethosu/te/dma.py
##########
@@ -67,63 +67,72 @@ def _pad(*indices):
return _pad
-def read_compute(tensor: te.Tensor, layout: str, zero_point: int, scale: float) -> te.Tensor:
+def read_compute(tensor: te.Tensor, zero_point: int, scale: float, layout: str = None) -> te.Tensor:
"""A tensor expression which represents a read.
Parameters
----------
tensor : te.Tensor
The tensor to read.
- layout : str
- The layout of the tensor, either NHWC or NHCWB16.
zero_point : int
The zero point of the tensor.
scale : float
The scale of the tensor.
+ layout : Optional[str]
+ The layout of the tensor, either NHWC or NHCWB16.
Returns
-------
te.Tensor
The tensor having been read.
"""
- assert layout in {"NHWC", "NHCWB16"}
read_attrs = {
"op": "ethosu_read",
- "layout": layout,
"zero_point": zero_point,
"scale": scale,
}
+
+ if layout:
+ assert layout in {"NHWC", "NHCWB16"}
+ read_attrs["layout"] = layout
+
return te.compute(tensor.shape, lambda *i: tensor(*i), name="ethosu_read", attrs=read_attrs)
-def write_compute(tensor: te.Tensor, layout: str, zero_point: int, scale: float) -> te.Tensor:
+def write_compute(
+ tensor: te.Tensor, zero_point: int, scale: float, layout: str = None
Review comment:
```suggestion
tensor: te.Tensor, zero_point: int, scale: float, layout: Optional[str] = None
```
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ekalda commented on a change in pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
Posted by GitBox <gi...@apache.org>.
ekalda commented on a change in pull request #9457:
URL: https://github.com/apache/tvm/pull/9457#discussion_r750118269
##########
File path: python/tvm/relay/backend/contrib/ethosu/op/identity.py
##########
@@ -0,0 +1,98 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+# pylint: disable=unused-argument
+"""Relay identity operator"""
+
+import tvm
+from tvm.relay.op import _make
+from tvm.topi.generic import schedule_injective
+from tvm.relay.op.op import OpStrategy
+from tvm.relay.op import strategy as _strategy
+
+from ..te import identity_compute
+
+
+@tvm.ir.register_op_attr("contrib.ethosu.identity", "FTVMCompute")
+def create_ethosu_identity_compute(attrs, args, out_type):
+ """Create an ethosu_identity compute op."""
+ ifm = args[0]
+ lut = args[1]
+ ifm_scale = attrs.ifm_scale
+ ifm_zero_point = attrs.ifm_zero_point
+ ofm_scale = attrs.ofm_scale
+ ofm_zero_point = attrs.ofm_zero_point
+ activation = attrs.activation
+ op = identity_compute(
+ ifm, lut, ifm_scale, ifm_zero_point, ofm_scale, ofm_zero_point, activation
+ )
+ return [op]
+
+
+@tvm.ir.register_op_attr("contrib.ethosu.identity", "FTVMStrategy")
+def pooling_strategy_ethosu(attrs, inputs, out_type, target):
Review comment:
Done
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ekalda commented on a change in pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
Posted by GitBox <gi...@apache.org>.
ekalda commented on a change in pull request #9457:
URL: https://github.com/apache/tvm/pull/9457#discussion_r750120346
##########
File path: tests/python/contrib/test_ethosu/test_type_inference.py
##########
@@ -171,5 +180,26 @@ def test_ethosu_pooling_invalid_dtype():
run_opt_pass(func, relay.transform.InferType())
+def test_ethosu_identity_invalid_shape():
+ invalid_shape = [1, 2, 3, 4, 5]
+ dtype = "int8"
+ ifm = relay.var("ifm", shape=invalid_shape, dtype=dtype)
+
+ identity = make_ethosu_identity(ifm)
+ func = relay.Function([ifm], identity)
+ with pytest.raises(TVMError):
+ run_opt_pass(func, relay.transform.InferType())
+
+
+def test_ethosu_invalid_invalid_dtype():
Review comment:
Oops, done
##########
File path: tests/python/contrib/test_ethosu/test_type_inference.py
##########
@@ -171,5 +180,26 @@ def test_ethosu_pooling_invalid_dtype():
run_opt_pass(func, relay.transform.InferType())
+def test_ethosu_identity_invalid_shape():
Review comment:
Done
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] mbaret commented on a change in pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
Posted by GitBox <gi...@apache.org>.
mbaret commented on a change in pull request #9457:
URL: https://github.com/apache/tvm/pull/9457#discussion_r743789916
##########
File path: python/tvm/relay/backend/contrib/ethosu/op/identity.py
##########
@@ -0,0 +1,98 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+# pylint: disable=unused-argument
+"""Relay identity operator"""
+
+import tvm
+from tvm.relay.op import _make
+from tvm.topi.generic import schedule_injective
+from tvm.relay.op.op import OpStrategy
+from tvm.relay.op import strategy as _strategy
+
+from ..te import identity_compute
+
+
+@tvm.ir.register_op_attr("contrib.ethosu.identity", "FTVMCompute")
+def create_ethosu_identity_compute(attrs, args, out_type):
+ """Create an ethosu_identity compute op."""
+ ifm = args[0]
+ lut = args[1]
+ ifm_scale = attrs.ifm_scale
+ ifm_zero_point = attrs.ifm_zero_point
+ ofm_scale = attrs.ofm_scale
+ ofm_zero_point = attrs.ofm_zero_point
+ activation = attrs.activation
+ op = identity_compute(
+ ifm, lut, ifm_scale, ifm_zero_point, ofm_scale, ofm_zero_point, activation
+ )
+ return [op]
+
+
+@tvm.ir.register_op_attr("contrib.ethosu.identity", "FTVMStrategy")
+def pooling_strategy_ethosu(attrs, inputs, out_type, target):
Review comment:
identity_strategy_ethosu
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir/identity.py
##########
@@ -0,0 +1,87 @@
+# 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
+"""Extract information from the identity operator in TIR."""
+from typing import Dict, Tuple
+import tvm
+from .dma import get_read_params, get_write_params
+from .spec import SerialKernel, SerialActivation, SerialPooling, SerialPadding
+from .utils import get_op_attrs
+
+
+def get_identity_params(
+ stmt: tvm.tir.AttrStmt,
+ producers: Dict[tvm.tir.Var, tvm.tir.AttrStmt],
+ consumers: Dict[tvm.tir.Var, tvm.tir.AttrStmt],
+) -> Tuple[SerialPooling, tvm.tir.Var, tvm.tir.Var]:
+ """Get the parameters necessary to construct a call_extern for a pooling.
+
+ Parameters
+ ----------
+ stmt : tvm.tir.AttrStmt
+ The outermost attribute statement of a convolution loop nest.
+ producers : Dict[tvm.tir.Var, tvm.tir.AttrStmt]
+ A dictionary to associate pointers with the loop nest
+ that produces their values.
+ consumers : Dict[tvm.tir.Var, tvm.tir.AttrStmt]
+ A dictionary to associate pointers with the loop nest
+ that consumes their values.
+
+ Returns
+ -------
+ SerialPooling
+ The parameters needed to construct a 2D pooling.
+ output_pointer : tvm.tir.Var
+ The output pointer of the pooling operation.
+ replace_pointer : tvm.tir.Var
+ The output pointer of the DMA write operation, which is to replace
+ the pooling output pointer.
+
+ """
+ attrs, _ = get_op_attrs(stmt)
+ # Find the inner loop
+ while hasattr(stmt, "body"):
+ stmt = stmt.body
+
+ input_pointer = stmt.value.buffer_var
+ output_pointer = stmt.buffer_var
+
+ read = producers[input_pointer]
+ write = consumers[output_pointer]
+
+ serial_ifm, _, _ = get_read_params(read)
+ serial_ofm, _, write_output_pointer = get_write_params(write)
+
+ replace_pointer = write_output_pointer
+
+ # TODO (maybe): Support stand alone RELU through clamping in identity
Review comment:
Clarify this
##########
File path: tests/python/contrib/test_ethosu/test_type_inference.py
##########
@@ -171,5 +180,26 @@ def test_ethosu_pooling_invalid_dtype():
run_opt_pass(func, relay.transform.InferType())
+def test_ethosu_identity_invalid_shape():
+ invalid_shape = [1, 2, 3, 4, 5]
+ dtype = "int8"
+ ifm = relay.var("ifm", shape=invalid_shape, dtype=dtype)
+
+ identity = make_ethosu_identity(ifm)
+ func = relay.Function([ifm], identity)
+ with pytest.raises(TVMError):
+ run_opt_pass(func, relay.transform.InferType())
+
+
+def test_ethosu_invalid_invalid_dtype():
Review comment:
test_ethosu_identity_invalid_dtype
##########
File path: python/tvm/relay/backend/contrib/ethosu/legalize.py
##########
@@ -423,11 +526,15 @@ class LegalizeEthosU:
def transform_module(
self, mod: tvm.ir.IRModule, ctx: tvm.ir.transform.PassContext
) -> tvm.ir.IRModule:
+ """Legalize the oerators that can be offloaded to the NPU"""
Review comment:
operators
##########
File path: tests/python/contrib/test_ethosu/test_type_inference.py
##########
@@ -171,5 +180,26 @@ def test_ethosu_pooling_invalid_dtype():
run_opt_pass(func, relay.transform.InferType())
+def test_ethosu_identity_invalid_shape():
Review comment:
I think we need to test the valid cases as well.
##########
File path: python/tvm/relay/op/contrib/ethosu.py
##########
@@ -481,6 +493,8 @@ def pattern_table() -> List[Tuple[str, tvm.relay.dataflow_pattern.DFPattern, Cal
qnn_avgpool2d_pattern(),
lambda pat: AvgPool2DParams(pat).is_valid(),
),
+ ("ethosu.strided_slice", strided_slice_pattern(), lambda pat: True),
Review comment:
Let's put the proper restrictions here so we don't accidentally offload, for example, floating point tensors.
##########
File path: src/relay/op/contrib/ethosu/identity.cc
##########
@@ -0,0 +1,128 @@
+/*
+ * 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/op/contrib/ethosu/identity.cc
+ * \brief Property def of the Arm Ethos-U identity op.
+ */
+#include <tvm/relay/op.h>
+
+#include "common.h"
+
+namespace tvm {
+namespace relay {
+namespace op {
+namespace contrib {
+namespace ethosu {
+
+/*! \brief Attributes used by the Ethos(TM)-U NPU identity operator */
+struct EthosuIdentityAttrs : public tvm::AttrsNode<EthosuIdentityAttrs> {
+ double ifm_scale;
+ int ifm_zero_point;
+ double ofm_scale;
+ int ofm_zero_point;
+ String activation;
+
+ TVM_DECLARE_ATTRS(EthosuIdentityAttrs, "relay.attrs.EthosuIdentityAttrs") {
+ TVM_ATTR_FIELD(ifm_scale).describe("The quantization scale for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ifm_zero_point)
+ .describe("The quantization zero point for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_scale).describe("The quantization scale for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_zero_point)
+ .describe("The quantization zero point for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(activation)
+ .describe(
+ "The activation function to use. "
+ "'NONE' - no activation function. "
+ "'TANH' - tanh activation function. "
+ "'SIGMOID' - sigmoid activation function. "
+ "'LUT' - use a look-up table to perform the activation function.")
+ .set_default("NONE");
+ }
+};
+
+TVM_REGISTER_NODE_TYPE(EthosuIdentityAttrs);
+
+bool EthosuIdentityRel(const Array<Type>& types, int num_inputs, const Attrs& attrs,
+ const TypeReporter& reporter) {
+ int ifm_index = 0;
+ int result_index = 2;
+ ICHECK_EQ(types.size(), result_index + 1);
+
+ const auto* ifm = types[ifm_index].as<TensorTypeNode>();
+ if (ifm == nullptr) return false;
+
+ const auto* param = attrs.as<EthosuIdentityAttrs>();
+
+ ICHECK(param != nullptr) << "EthosuIdentityAttrs cannot be nullptr.";
+
+ if (ifm->dtype != DataType::UInt(8) && ifm->dtype != DataType::Int(8)) {
+ reporter->GetDiagCtx().EmitFatal(
+ Diagnostic::Error(reporter->GetSpan())
+ << "Invalid operator: Expected type(uint8) or type(int8) for ifm but was " << ifm->dtype);
+ return false;
+ }
+
+ if (ifm->shape.size() > 4) {
+ reporter->GetDiagCtx().EmitFatal(
+ Diagnostic::Error(reporter->GetSpan())
+ << "Invalid operator: Input Feature Map should be at most 4 dimensional, but was "
+ << ifm->shape);
+ return false;
+ }
+
+ // Assign ofm type
+ auto ofm_shape = ifm->shape;
+ reporter->Assign(types[result_index], TensorType(ofm_shape, ifm->dtype));
+ return true;
+}
+
+Expr MakeEthosuIdentity(Expr ifm, Expr lut, double ifm_scale, int ifm_zero_point, double ofm_scale,
+ int ofm_zero_point, String activation) {
+ auto attrs = make_object<EthosuIdentityAttrs>();
+ attrs->ifm_scale = ifm_scale;
+ attrs->ifm_zero_point = ifm_zero_point;
+ attrs->ofm_scale = ofm_scale;
+ attrs->ofm_zero_point = ofm_zero_point;
+ attrs->activation = std::move(activation);
+ static const Op& op = Op::Get("contrib.ethosu.identity");
+ return Call(op, {ifm, lut}, Attrs(attrs), {});
+}
+
+TVM_REGISTER_GLOBAL("relay.op._make.ethosu_identity").set_body_typed(MakeEthosuIdentity);
+
+RELAY_REGISTER_OP("contrib.ethosu.identity")
+ .describe(R"code(Identity operator for Ethos-U NPUs.
+
+This Relay operator performs the identity operation on Ethos(TM)-U NPU with a capability
+to requantize the data. It accepts input with any shape that is less or equal to 4.
Review comment:
Maybe 'It accepts input tensors of 4 dimensions or less.'?
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir/identity.py
##########
@@ -0,0 +1,87 @@
+# 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
+"""Extract information from the identity operator in TIR."""
+from typing import Dict, Tuple
+import tvm
+from .dma import get_read_params, get_write_params
+from .spec import SerialKernel, SerialActivation, SerialPooling, SerialPadding
+from .utils import get_op_attrs
+
+
+def get_identity_params(
+ stmt: tvm.tir.AttrStmt,
+ producers: Dict[tvm.tir.Var, tvm.tir.AttrStmt],
+ consumers: Dict[tvm.tir.Var, tvm.tir.AttrStmt],
+) -> Tuple[SerialPooling, tvm.tir.Var, tvm.tir.Var]:
+ """Get the parameters necessary to construct a call_extern for a pooling.
+
+ Parameters
+ ----------
+ stmt : tvm.tir.AttrStmt
+ The outermost attribute statement of a convolution loop nest.
Review comment:
identity pooling loop nest
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir/identity.py
##########
@@ -0,0 +1,87 @@
+# 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
+"""Extract information from the identity operator in TIR."""
+from typing import Dict, Tuple
+import tvm
+from .dma import get_read_params, get_write_params
+from .spec import SerialKernel, SerialActivation, SerialPooling, SerialPadding
+from .utils import get_op_attrs
+
+
+def get_identity_params(
+ stmt: tvm.tir.AttrStmt,
+ producers: Dict[tvm.tir.Var, tvm.tir.AttrStmt],
+ consumers: Dict[tvm.tir.Var, tvm.tir.AttrStmt],
+) -> Tuple[SerialPooling, tvm.tir.Var, tvm.tir.Var]:
+ """Get the parameters necessary to construct a call_extern for a pooling.
Review comment:
for an identity pooling
##########
File path: tests/python/contrib/test_ethosu/test_codegen.py
##########
@@ -343,5 +343,178 @@ def representative_dataset():
infra.verify_source(compiled_models, accel_type)
+@pytest.mark.parametrize("accel_type", ACCEL_TYPES)
+@pytest.mark.parametrize("ifm_shape", [(3, 2), (1, 15, 11, 7), (3, 1, 12), (400,)])
+@pytest.mark.parametrize("ifm_scale, ifm_zp, ofm_scale, ofm_zp", [(1, 0, 1, 0), (0.015, 3, 0.2, 5)])
+def test_ethosu_identity_codegen(ifm_shape, ifm_scale, ifm_zp, ofm_scale, ofm_zp, accel_type):
+ # Create a "partitioned" Relay function
+ ifm = relay.var("ifm", shape=ifm_shape, dtype="int8")
+ ifm0 = relay.var("ifm0", shape=ifm_shape, dtype="int8")
+ identity = infra.make_ethosu_identity(
+ ifm0, ifm_scale=ifm_scale, ifm_zero_point=ifm_zp, ofm_scale=ofm_scale, ofm_zero_point=ofm_zp
+ )
+ glb_ethosu = relay.GlobalVar("tvmgen_default_ethosu_main_0")
+
+ func = (
+ relay.Function([ifm0], identity)
+ .with_attr("Inline", 1)
+ .with_attr("Compiler", "ethosu")
+ .with_attr("global_symbol", "tvmgen_default_ethosu_main_0")
+ .with_attr("Primitive", 1)
+ )
+ mod = tvm.IRModule()
+ mod[glb_ethosu] = func
+ mod = relay.transform.InferType()(mod)
+
+ call = relay.Call(glb_ethosu, [ifm])
+ mod["main"] = relay.Function([ifm], call)
+ mod = relay.transform.InferType()(mod)
+
+ in_data = np.random.randint(-120, high=120, size=ifm_shape, dtype="int8")
+ requant_data = (ifm_scale * (in_data - ifm_zp)) / ofm_scale + ofm_zp
+ out_data = np.round(np.clip(requant_data, -128, 127)).astype("int8")
+
+ compiled_model = infra.build_source(
+ mod, {"ifm": in_data}, [out_data], accel_type, output_tolerance=1
+ )
+
+ imported_modules = compiled_model[0].executor_factory.lib.imported_modules
+ assert len(imported_modules) == 2
+ ethosu_module = imported_modules[0]
+
+ # Verify generated C source
+ get_cs = tvm._ffi.get_global_func("runtime.module.ethosu.getcs")
+ cmms = get_cs(ethosu_module)
+ cmms = bytes.fromhex(cmms)
+
+ infra.print_payload(cmms)
+ infra.verify_source(compiled_model, accel_type)
+
+
+@pytest.mark.parametrize("accel_type", ACCEL_TYPES)
+@pytest.mark.parametrize(
+ "ifm_shape, new_shape",
+ [
+ ((1, 4, 1, 2), (1, 1, 1, 8)),
+ (
+ (
+ 5,
+ 1,
+ 20,
+ ),
+ (1, 5, 1, 20),
+ ),
+ ((12, 20), (1, 6, 4, 10)),
+ ((12, 20), (6, 4, 10)),
+ ((20,), (4, 5)),
+ ],
+)
+def test_relay_reshape_codegen(ifm_shape, new_shape, accel_type):
+ # Create a "partitioned" Relay graph
Review comment:
It seems like some of this could be refactored out into a common function and reused between tests.
##########
File path: tests/python/contrib/test_ethosu/test_codegen.py
##########
@@ -343,5 +343,178 @@ def representative_dataset():
infra.verify_source(compiled_models, accel_type)
+@pytest.mark.parametrize("accel_type", ACCEL_TYPES)
+@pytest.mark.parametrize("ifm_shape", [(3, 2), (1, 15, 11, 7), (3, 1, 12), (400,)])
+@pytest.mark.parametrize("ifm_scale, ifm_zp, ofm_scale, ofm_zp", [(1, 0, 1, 0), (0.015, 3, 0.2, 5)])
+def test_ethosu_identity_codegen(ifm_shape, ifm_scale, ifm_zp, ofm_scale, ofm_zp, accel_type):
+ # Create a "partitioned" Relay function
+ ifm = relay.var("ifm", shape=ifm_shape, dtype="int8")
+ ifm0 = relay.var("ifm0", shape=ifm_shape, dtype="int8")
+ identity = infra.make_ethosu_identity(
+ ifm0, ifm_scale=ifm_scale, ifm_zero_point=ifm_zp, ofm_scale=ofm_scale, ofm_zero_point=ofm_zp
+ )
+ glb_ethosu = relay.GlobalVar("tvmgen_default_ethosu_main_0")
+
+ func = (
+ relay.Function([ifm0], identity)
+ .with_attr("Inline", 1)
+ .with_attr("Compiler", "ethosu")
+ .with_attr("global_symbol", "tvmgen_default_ethosu_main_0")
+ .with_attr("Primitive", 1)
+ )
+ mod = tvm.IRModule()
+ mod[glb_ethosu] = func
+ mod = relay.transform.InferType()(mod)
+
+ call = relay.Call(glb_ethosu, [ifm])
+ mod["main"] = relay.Function([ifm], call)
+ mod = relay.transform.InferType()(mod)
+
+ in_data = np.random.randint(-120, high=120, size=ifm_shape, dtype="int8")
+ requant_data = (ifm_scale * (in_data - ifm_zp)) / ofm_scale + ofm_zp
+ out_data = np.round(np.clip(requant_data, -128, 127)).astype("int8")
+
+ compiled_model = infra.build_source(
+ mod, {"ifm": in_data}, [out_data], accel_type, output_tolerance=1
+ )
+
+ imported_modules = compiled_model[0].executor_factory.lib.imported_modules
+ assert len(imported_modules) == 2
+ ethosu_module = imported_modules[0]
+
+ # Verify generated C source
+ get_cs = tvm._ffi.get_global_func("runtime.module.ethosu.getcs")
+ cmms = get_cs(ethosu_module)
+ cmms = bytes.fromhex(cmms)
+
+ infra.print_payload(cmms)
+ infra.verify_source(compiled_model, accel_type)
+
+
+@pytest.mark.parametrize("accel_type", ACCEL_TYPES)
+@pytest.mark.parametrize(
Review comment:
We should either test the 'special' indices for reshape (-1 and -2), or we should explicitly disallow these.
##########
File path: python/tvm/relay/backend/contrib/ethosu/op/identity.py
##########
@@ -0,0 +1,98 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+# pylint: disable=unused-argument
+"""Relay identity operator"""
+
+import tvm
+from tvm.relay.op import _make
+from tvm.topi.generic import schedule_injective
+from tvm.relay.op.op import OpStrategy
+from tvm.relay.op import strategy as _strategy
+
+from ..te import identity_compute
+
+
+@tvm.ir.register_op_attr("contrib.ethosu.identity", "FTVMCompute")
+def create_ethosu_identity_compute(attrs, args, out_type):
+ """Create an ethosu_identity compute op."""
+ ifm = args[0]
+ lut = args[1]
+ ifm_scale = attrs.ifm_scale
+ ifm_zero_point = attrs.ifm_zero_point
+ ofm_scale = attrs.ofm_scale
+ ofm_zero_point = attrs.ofm_zero_point
+ activation = attrs.activation
+ op = identity_compute(
+ ifm, lut, ifm_scale, ifm_zero_point, ofm_scale, ofm_zero_point, activation
+ )
+ return [op]
+
+
+@tvm.ir.register_op_attr("contrib.ethosu.identity", "FTVMStrategy")
+def pooling_strategy_ethosu(attrs, inputs, out_type, target):
Review comment:
identity_strategy_ethosu
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir/identity.py
##########
@@ -0,0 +1,87 @@
+# 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
+"""Extract information from the identity operator in TIR."""
+from typing import Dict, Tuple
+import tvm
+from .dma import get_read_params, get_write_params
+from .spec import SerialKernel, SerialActivation, SerialPooling, SerialPadding
+from .utils import get_op_attrs
+
+
+def get_identity_params(
+ stmt: tvm.tir.AttrStmt,
+ producers: Dict[tvm.tir.Var, tvm.tir.AttrStmt],
+ consumers: Dict[tvm.tir.Var, tvm.tir.AttrStmt],
+) -> Tuple[SerialPooling, tvm.tir.Var, tvm.tir.Var]:
+ """Get the parameters necessary to construct a call_extern for a pooling.
+
+ Parameters
+ ----------
+ stmt : tvm.tir.AttrStmt
+ The outermost attribute statement of a convolution loop nest.
+ producers : Dict[tvm.tir.Var, tvm.tir.AttrStmt]
+ A dictionary to associate pointers with the loop nest
+ that produces their values.
+ consumers : Dict[tvm.tir.Var, tvm.tir.AttrStmt]
+ A dictionary to associate pointers with the loop nest
+ that consumes their values.
+
+ Returns
+ -------
+ SerialPooling
+ The parameters needed to construct a 2D pooling.
+ output_pointer : tvm.tir.Var
+ The output pointer of the pooling operation.
+ replace_pointer : tvm.tir.Var
+ The output pointer of the DMA write operation, which is to replace
+ the pooling output pointer.
+
+ """
+ attrs, _ = get_op_attrs(stmt)
+ # Find the inner loop
+ while hasattr(stmt, "body"):
+ stmt = stmt.body
+
+ input_pointer = stmt.value.buffer_var
+ output_pointer = stmt.buffer_var
+
+ read = producers[input_pointer]
+ write = consumers[output_pointer]
+
+ serial_ifm, _, _ = get_read_params(read)
+ serial_ofm, _, write_output_pointer = get_write_params(write)
+
+ replace_pointer = write_output_pointer
+
+ # TODO (maybe): Support stand alone RELU through clamping in identity
Review comment:
Clarify this
##########
File path: tests/python/contrib/test_ethosu/test_type_inference.py
##########
@@ -171,5 +180,26 @@ def test_ethosu_pooling_invalid_dtype():
run_opt_pass(func, relay.transform.InferType())
+def test_ethosu_identity_invalid_shape():
+ invalid_shape = [1, 2, 3, 4, 5]
+ dtype = "int8"
+ ifm = relay.var("ifm", shape=invalid_shape, dtype=dtype)
+
+ identity = make_ethosu_identity(ifm)
+ func = relay.Function([ifm], identity)
+ with pytest.raises(TVMError):
+ run_opt_pass(func, relay.transform.InferType())
+
+
+def test_ethosu_invalid_invalid_dtype():
Review comment:
test_ethosu_identity_invalid_dtype
##########
File path: python/tvm/relay/backend/contrib/ethosu/legalize.py
##########
@@ -423,11 +526,15 @@ class LegalizeEthosU:
def transform_module(
self, mod: tvm.ir.IRModule, ctx: tvm.ir.transform.PassContext
) -> tvm.ir.IRModule:
+ """Legalize the oerators that can be offloaded to the NPU"""
Review comment:
operators
##########
File path: tests/python/contrib/test_ethosu/test_type_inference.py
##########
@@ -171,5 +180,26 @@ def test_ethosu_pooling_invalid_dtype():
run_opt_pass(func, relay.transform.InferType())
+def test_ethosu_identity_invalid_shape():
Review comment:
I think we need to test the valid cases as well.
##########
File path: python/tvm/relay/op/contrib/ethosu.py
##########
@@ -481,6 +493,8 @@ def pattern_table() -> List[Tuple[str, tvm.relay.dataflow_pattern.DFPattern, Cal
qnn_avgpool2d_pattern(),
lambda pat: AvgPool2DParams(pat).is_valid(),
),
+ ("ethosu.strided_slice", strided_slice_pattern(), lambda pat: True),
Review comment:
Let's put the proper restrictions here so we don't accidentally offload, for example, floating point tensors.
##########
File path: src/relay/op/contrib/ethosu/identity.cc
##########
@@ -0,0 +1,128 @@
+/*
+ * 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/op/contrib/ethosu/identity.cc
+ * \brief Property def of the Arm Ethos-U identity op.
+ */
+#include <tvm/relay/op.h>
+
+#include "common.h"
+
+namespace tvm {
+namespace relay {
+namespace op {
+namespace contrib {
+namespace ethosu {
+
+/*! \brief Attributes used by the Ethos(TM)-U NPU identity operator */
+struct EthosuIdentityAttrs : public tvm::AttrsNode<EthosuIdentityAttrs> {
+ double ifm_scale;
+ int ifm_zero_point;
+ double ofm_scale;
+ int ofm_zero_point;
+ String activation;
+
+ TVM_DECLARE_ATTRS(EthosuIdentityAttrs, "relay.attrs.EthosuIdentityAttrs") {
+ TVM_ATTR_FIELD(ifm_scale).describe("The quantization scale for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ifm_zero_point)
+ .describe("The quantization zero point for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_scale).describe("The quantization scale for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_zero_point)
+ .describe("The quantization zero point for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(activation)
+ .describe(
+ "The activation function to use. "
+ "'NONE' - no activation function. "
+ "'TANH' - tanh activation function. "
+ "'SIGMOID' - sigmoid activation function. "
+ "'LUT' - use a look-up table to perform the activation function.")
+ .set_default("NONE");
+ }
+};
+
+TVM_REGISTER_NODE_TYPE(EthosuIdentityAttrs);
+
+bool EthosuIdentityRel(const Array<Type>& types, int num_inputs, const Attrs& attrs,
+ const TypeReporter& reporter) {
+ int ifm_index = 0;
+ int result_index = 2;
+ ICHECK_EQ(types.size(), result_index + 1);
+
+ const auto* ifm = types[ifm_index].as<TensorTypeNode>();
+ if (ifm == nullptr) return false;
+
+ const auto* param = attrs.as<EthosuIdentityAttrs>();
+
+ ICHECK(param != nullptr) << "EthosuIdentityAttrs cannot be nullptr.";
+
+ if (ifm->dtype != DataType::UInt(8) && ifm->dtype != DataType::Int(8)) {
+ reporter->GetDiagCtx().EmitFatal(
+ Diagnostic::Error(reporter->GetSpan())
+ << "Invalid operator: Expected type(uint8) or type(int8) for ifm but was " << ifm->dtype);
+ return false;
+ }
+
+ if (ifm->shape.size() > 4) {
+ reporter->GetDiagCtx().EmitFatal(
+ Diagnostic::Error(reporter->GetSpan())
+ << "Invalid operator: Input Feature Map should be at most 4 dimensional, but was "
+ << ifm->shape);
+ return false;
+ }
+
+ // Assign ofm type
+ auto ofm_shape = ifm->shape;
+ reporter->Assign(types[result_index], TensorType(ofm_shape, ifm->dtype));
+ return true;
+}
+
+Expr MakeEthosuIdentity(Expr ifm, Expr lut, double ifm_scale, int ifm_zero_point, double ofm_scale,
+ int ofm_zero_point, String activation) {
+ auto attrs = make_object<EthosuIdentityAttrs>();
+ attrs->ifm_scale = ifm_scale;
+ attrs->ifm_zero_point = ifm_zero_point;
+ attrs->ofm_scale = ofm_scale;
+ attrs->ofm_zero_point = ofm_zero_point;
+ attrs->activation = std::move(activation);
+ static const Op& op = Op::Get("contrib.ethosu.identity");
+ return Call(op, {ifm, lut}, Attrs(attrs), {});
+}
+
+TVM_REGISTER_GLOBAL("relay.op._make.ethosu_identity").set_body_typed(MakeEthosuIdentity);
+
+RELAY_REGISTER_OP("contrib.ethosu.identity")
+ .describe(R"code(Identity operator for Ethos-U NPUs.
+
+This Relay operator performs the identity operation on Ethos(TM)-U NPU with a capability
+to requantize the data. It accepts input with any shape that is less or equal to 4.
Review comment:
Maybe 'It accepts input tensors of 4 dimensions or less.'?
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir/identity.py
##########
@@ -0,0 +1,87 @@
+# 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
+"""Extract information from the identity operator in TIR."""
+from typing import Dict, Tuple
+import tvm
+from .dma import get_read_params, get_write_params
+from .spec import SerialKernel, SerialActivation, SerialPooling, SerialPadding
+from .utils import get_op_attrs
+
+
+def get_identity_params(
+ stmt: tvm.tir.AttrStmt,
+ producers: Dict[tvm.tir.Var, tvm.tir.AttrStmt],
+ consumers: Dict[tvm.tir.Var, tvm.tir.AttrStmt],
+) -> Tuple[SerialPooling, tvm.tir.Var, tvm.tir.Var]:
+ """Get the parameters necessary to construct a call_extern for a pooling.
+
+ Parameters
+ ----------
+ stmt : tvm.tir.AttrStmt
+ The outermost attribute statement of a convolution loop nest.
Review comment:
identity pooling loop nest
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir/identity.py
##########
@@ -0,0 +1,87 @@
+# 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
+"""Extract information from the identity operator in TIR."""
+from typing import Dict, Tuple
+import tvm
+from .dma import get_read_params, get_write_params
+from .spec import SerialKernel, SerialActivation, SerialPooling, SerialPadding
+from .utils import get_op_attrs
+
+
+def get_identity_params(
+ stmt: tvm.tir.AttrStmt,
+ producers: Dict[tvm.tir.Var, tvm.tir.AttrStmt],
+ consumers: Dict[tvm.tir.Var, tvm.tir.AttrStmt],
+) -> Tuple[SerialPooling, tvm.tir.Var, tvm.tir.Var]:
+ """Get the parameters necessary to construct a call_extern for a pooling.
Review comment:
for an identity pooling
##########
File path: tests/python/contrib/test_ethosu/test_codegen.py
##########
@@ -343,5 +343,178 @@ def representative_dataset():
infra.verify_source(compiled_models, accel_type)
+@pytest.mark.parametrize("accel_type", ACCEL_TYPES)
+@pytest.mark.parametrize("ifm_shape", [(3, 2), (1, 15, 11, 7), (3, 1, 12), (400,)])
+@pytest.mark.parametrize("ifm_scale, ifm_zp, ofm_scale, ofm_zp", [(1, 0, 1, 0), (0.015, 3, 0.2, 5)])
+def test_ethosu_identity_codegen(ifm_shape, ifm_scale, ifm_zp, ofm_scale, ofm_zp, accel_type):
+ # Create a "partitioned" Relay function
+ ifm = relay.var("ifm", shape=ifm_shape, dtype="int8")
+ ifm0 = relay.var("ifm0", shape=ifm_shape, dtype="int8")
+ identity = infra.make_ethosu_identity(
+ ifm0, ifm_scale=ifm_scale, ifm_zero_point=ifm_zp, ofm_scale=ofm_scale, ofm_zero_point=ofm_zp
+ )
+ glb_ethosu = relay.GlobalVar("tvmgen_default_ethosu_main_0")
+
+ func = (
+ relay.Function([ifm0], identity)
+ .with_attr("Inline", 1)
+ .with_attr("Compiler", "ethosu")
+ .with_attr("global_symbol", "tvmgen_default_ethosu_main_0")
+ .with_attr("Primitive", 1)
+ )
+ mod = tvm.IRModule()
+ mod[glb_ethosu] = func
+ mod = relay.transform.InferType()(mod)
+
+ call = relay.Call(glb_ethosu, [ifm])
+ mod["main"] = relay.Function([ifm], call)
+ mod = relay.transform.InferType()(mod)
+
+ in_data = np.random.randint(-120, high=120, size=ifm_shape, dtype="int8")
+ requant_data = (ifm_scale * (in_data - ifm_zp)) / ofm_scale + ofm_zp
+ out_data = np.round(np.clip(requant_data, -128, 127)).astype("int8")
+
+ compiled_model = infra.build_source(
+ mod, {"ifm": in_data}, [out_data], accel_type, output_tolerance=1
+ )
+
+ imported_modules = compiled_model[0].executor_factory.lib.imported_modules
+ assert len(imported_modules) == 2
+ ethosu_module = imported_modules[0]
+
+ # Verify generated C source
+ get_cs = tvm._ffi.get_global_func("runtime.module.ethosu.getcs")
+ cmms = get_cs(ethosu_module)
+ cmms = bytes.fromhex(cmms)
+
+ infra.print_payload(cmms)
+ infra.verify_source(compiled_model, accel_type)
+
+
+@pytest.mark.parametrize("accel_type", ACCEL_TYPES)
+@pytest.mark.parametrize(
+ "ifm_shape, new_shape",
+ [
+ ((1, 4, 1, 2), (1, 1, 1, 8)),
+ (
+ (
+ 5,
+ 1,
+ 20,
+ ),
+ (1, 5, 1, 20),
+ ),
+ ((12, 20), (1, 6, 4, 10)),
+ ((12, 20), (6, 4, 10)),
+ ((20,), (4, 5)),
+ ],
+)
+def test_relay_reshape_codegen(ifm_shape, new_shape, accel_type):
+ # Create a "partitioned" Relay graph
Review comment:
It seems like some of this could be refactored out into a common function and reused between tests.
##########
File path: tests/python/contrib/test_ethosu/test_codegen.py
##########
@@ -343,5 +343,178 @@ def representative_dataset():
infra.verify_source(compiled_models, accel_type)
+@pytest.mark.parametrize("accel_type", ACCEL_TYPES)
+@pytest.mark.parametrize("ifm_shape", [(3, 2), (1, 15, 11, 7), (3, 1, 12), (400,)])
+@pytest.mark.parametrize("ifm_scale, ifm_zp, ofm_scale, ofm_zp", [(1, 0, 1, 0), (0.015, 3, 0.2, 5)])
+def test_ethosu_identity_codegen(ifm_shape, ifm_scale, ifm_zp, ofm_scale, ofm_zp, accel_type):
+ # Create a "partitioned" Relay function
+ ifm = relay.var("ifm", shape=ifm_shape, dtype="int8")
+ ifm0 = relay.var("ifm0", shape=ifm_shape, dtype="int8")
+ identity = infra.make_ethosu_identity(
+ ifm0, ifm_scale=ifm_scale, ifm_zero_point=ifm_zp, ofm_scale=ofm_scale, ofm_zero_point=ofm_zp
+ )
+ glb_ethosu = relay.GlobalVar("tvmgen_default_ethosu_main_0")
+
+ func = (
+ relay.Function([ifm0], identity)
+ .with_attr("Inline", 1)
+ .with_attr("Compiler", "ethosu")
+ .with_attr("global_symbol", "tvmgen_default_ethosu_main_0")
+ .with_attr("Primitive", 1)
+ )
+ mod = tvm.IRModule()
+ mod[glb_ethosu] = func
+ mod = relay.transform.InferType()(mod)
+
+ call = relay.Call(glb_ethosu, [ifm])
+ mod["main"] = relay.Function([ifm], call)
+ mod = relay.transform.InferType()(mod)
+
+ in_data = np.random.randint(-120, high=120, size=ifm_shape, dtype="int8")
+ requant_data = (ifm_scale * (in_data - ifm_zp)) / ofm_scale + ofm_zp
+ out_data = np.round(np.clip(requant_data, -128, 127)).astype("int8")
+
+ compiled_model = infra.build_source(
+ mod, {"ifm": in_data}, [out_data], accel_type, output_tolerance=1
+ )
+
+ imported_modules = compiled_model[0].executor_factory.lib.imported_modules
+ assert len(imported_modules) == 2
+ ethosu_module = imported_modules[0]
+
+ # Verify generated C source
+ get_cs = tvm._ffi.get_global_func("runtime.module.ethosu.getcs")
+ cmms = get_cs(ethosu_module)
+ cmms = bytes.fromhex(cmms)
+
+ infra.print_payload(cmms)
+ infra.verify_source(compiled_model, accel_type)
+
+
+@pytest.mark.parametrize("accel_type", ACCEL_TYPES)
+@pytest.mark.parametrize(
Review comment:
We should either test the 'special' indices for reshape (-1 and -2), or we should explicitly disallow these.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] lhutton1 commented on a change in pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
Posted by GitBox <gi...@apache.org>.
lhutton1 commented on a change in pull request #9457:
URL: https://github.com/apache/tvm/pull/9457#discussion_r744634984
##########
File path: src/relay/op/contrib/ethosu/identity.cc
##########
@@ -0,0 +1,128 @@
+/*
+ * 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/op/contrib/ethosu/identity.cc
+ * \brief Property def of the Arm Ethos-U identity op.
+ */
+#include <tvm/relay/op.h>
+
+#include "common.h"
+
+namespace tvm {
+namespace relay {
+namespace op {
+namespace contrib {
+namespace ethosu {
+
+/*! \brief Attributes used by the Ethos(TM)-U NPU identity operator */
+struct EthosuIdentityAttrs : public tvm::AttrsNode<EthosuIdentityAttrs> {
+ double ifm_scale;
+ int ifm_zero_point;
+ double ofm_scale;
+ int ofm_zero_point;
+ String activation;
+
+ TVM_DECLARE_ATTRS(EthosuIdentityAttrs, "relay.attrs.EthosuIdentityAttrs") {
+ TVM_ATTR_FIELD(ifm_scale).describe("The quantization scale for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ifm_zero_point)
+ .describe("The quantization zero point for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_scale).describe("The quantization scale for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_zero_point)
+ .describe("The quantization zero point for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(activation)
+ .describe(
+ "The activation function to use. "
+ "'NONE' - no activation function. "
+ "'TANH' - tanh activation function. "
+ "'SIGMOID' - sigmoid activation function. "
+ "'LUT' - use a look-up table to perform the activation function.")
+ .set_default("NONE");
+ }
+};
+
+TVM_REGISTER_NODE_TYPE(EthosuIdentityAttrs);
+
+bool EthosuIdentityRel(const Array<Type>& types, int num_inputs, const Attrs& attrs,
+ const TypeReporter& reporter) {
+ int ifm_index = 0;
+ int result_index = 2;
+ ICHECK_EQ(types.size(), result_index + 1);
+
+ const auto* ifm = types[ifm_index].as<TensorTypeNode>();
+ if (ifm == nullptr) return false;
+
+ const auto* param = attrs.as<EthosuIdentityAttrs>();
+
+ ICHECK(param != nullptr) << "EthosuIdentityAttrs cannot be nullptr.";
+
+ if (ifm->dtype != DataType::UInt(8) && ifm->dtype != DataType::Int(8)) {
+ reporter->GetDiagCtx().EmitFatal(
+ Diagnostic::Error(reporter->GetSpan())
+ << "Invalid operator: Expected type(uint8) or type(int8) for ifm but was " << ifm->dtype);
+ return false;
+ }
+
+ if (ifm->shape.size() > 4) {
+ reporter->GetDiagCtx().EmitFatal(
+ Diagnostic::Error(reporter->GetSpan())
+ << "Invalid operator: Input Feature Map should be at most 4 dimensional, but was "
+ << ifm->shape);
+ return false;
+ }
+
+ // Assign ofm type
+ auto ofm_shape = ifm->shape;
+ reporter->Assign(types[result_index], TensorType(ofm_shape, ifm->dtype));
+ return true;
+}
+
+Expr MakeEthosuIdentity(Expr ifm, Expr lut, double ifm_scale, int ifm_zero_point, double ofm_scale,
+ int ofm_zero_point, String activation) {
+ auto attrs = make_object<EthosuIdentityAttrs>();
+ attrs->ifm_scale = ifm_scale;
+ attrs->ifm_zero_point = ifm_zero_point;
+ attrs->ofm_scale = ofm_scale;
+ attrs->ofm_zero_point = ofm_zero_point;
+ attrs->activation = std::move(activation);
+ static const Op& op = Op::Get("contrib.ethosu.identity");
+ return Call(op, {ifm, lut}, Attrs(attrs), {});
+}
+
+TVM_REGISTER_GLOBAL("relay.op._make.ethosu_identity").set_body_typed(MakeEthosuIdentity);
+
+RELAY_REGISTER_OP("contrib.ethosu.identity")
+ .describe(R"code(Identity operator for Ethos-U NPUs.
Review comment:
```suggestion
.describe(R"code(Identity operator for Ethos(TM)-U NPUs.
```
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] mbaret commented on pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
Posted by GitBox <gi...@apache.org>.
mbaret commented on pull request #9457:
URL: https://github.com/apache/tvm/pull/9457#issuecomment-970496003
Thanks everyone, this is now merged!
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ekalda commented on a change in pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
Posted by GitBox <gi...@apache.org>.
ekalda commented on a change in pull request #9457:
URL: https://github.com/apache/tvm/pull/9457#discussion_r750119977
##########
File path: tests/python/contrib/test_ethosu/test_codegen.py
##########
@@ -343,5 +343,178 @@ def representative_dataset():
infra.verify_source(compiled_models, accel_type)
+@pytest.mark.parametrize("accel_type", ACCEL_TYPES)
+@pytest.mark.parametrize("ifm_shape", [(3, 2), (1, 15, 11, 7), (3, 1, 12), (400,)])
+@pytest.mark.parametrize("ifm_scale, ifm_zp, ofm_scale, ofm_zp", [(1, 0, 1, 0), (0.015, 3, 0.2, 5)])
+def test_ethosu_identity_codegen(ifm_shape, ifm_scale, ifm_zp, ofm_scale, ofm_zp, accel_type):
+ # Create a "partitioned" Relay function
+ ifm = relay.var("ifm", shape=ifm_shape, dtype="int8")
+ ifm0 = relay.var("ifm0", shape=ifm_shape, dtype="int8")
+ identity = infra.make_ethosu_identity(
+ ifm0, ifm_scale=ifm_scale, ifm_zero_point=ifm_zp, ofm_scale=ofm_scale, ofm_zero_point=ofm_zp
+ )
+ glb_ethosu = relay.GlobalVar("tvmgen_default_ethosu_main_0")
+
+ func = (
+ relay.Function([ifm0], identity)
+ .with_attr("Inline", 1)
+ .with_attr("Compiler", "ethosu")
+ .with_attr("global_symbol", "tvmgen_default_ethosu_main_0")
+ .with_attr("Primitive", 1)
+ )
+ mod = tvm.IRModule()
+ mod[glb_ethosu] = func
+ mod = relay.transform.InferType()(mod)
+
+ call = relay.Call(glb_ethosu, [ifm])
+ mod["main"] = relay.Function([ifm], call)
+ mod = relay.transform.InferType()(mod)
+
+ in_data = np.random.randint(-120, high=120, size=ifm_shape, dtype="int8")
+ requant_data = (ifm_scale * (in_data - ifm_zp)) / ofm_scale + ofm_zp
+ out_data = np.round(np.clip(requant_data, -128, 127)).astype("int8")
+
+ compiled_model = infra.build_source(
+ mod, {"ifm": in_data}, [out_data], accel_type, output_tolerance=1
+ )
+
+ imported_modules = compiled_model[0].executor_factory.lib.imported_modules
+ assert len(imported_modules) == 2
+ ethosu_module = imported_modules[0]
+
+ # Verify generated C source
+ get_cs = tvm._ffi.get_global_func("runtime.module.ethosu.getcs")
+ cmms = get_cs(ethosu_module)
+ cmms = bytes.fromhex(cmms)
+
+ infra.print_payload(cmms)
+ infra.verify_source(compiled_model, accel_type)
+
+
+@pytest.mark.parametrize("accel_type", ACCEL_TYPES)
+@pytest.mark.parametrize(
Review comment:
Done
##########
File path: tests/python/contrib/test_ethosu/test_codegen.py
##########
@@ -343,5 +343,178 @@ def representative_dataset():
infra.verify_source(compiled_models, accel_type)
+@pytest.mark.parametrize("accel_type", ACCEL_TYPES)
+@pytest.mark.parametrize("ifm_shape", [(3, 2), (1, 15, 11, 7), (3, 1, 12), (400,)])
+@pytest.mark.parametrize("ifm_scale, ifm_zp, ofm_scale, ofm_zp", [(1, 0, 1, 0), (0.015, 3, 0.2, 5)])
+def test_ethosu_identity_codegen(ifm_shape, ifm_scale, ifm_zp, ofm_scale, ofm_zp, accel_type):
+ # Create a "partitioned" Relay function
+ ifm = relay.var("ifm", shape=ifm_shape, dtype="int8")
+ ifm0 = relay.var("ifm0", shape=ifm_shape, dtype="int8")
+ identity = infra.make_ethosu_identity(
+ ifm0, ifm_scale=ifm_scale, ifm_zero_point=ifm_zp, ofm_scale=ofm_scale, ofm_zero_point=ofm_zp
+ )
+ glb_ethosu = relay.GlobalVar("tvmgen_default_ethosu_main_0")
+
+ func = (
+ relay.Function([ifm0], identity)
+ .with_attr("Inline", 1)
+ .with_attr("Compiler", "ethosu")
+ .with_attr("global_symbol", "tvmgen_default_ethosu_main_0")
+ .with_attr("Primitive", 1)
+ )
+ mod = tvm.IRModule()
+ mod[glb_ethosu] = func
+ mod = relay.transform.InferType()(mod)
+
+ call = relay.Call(glb_ethosu, [ifm])
+ mod["main"] = relay.Function([ifm], call)
+ mod = relay.transform.InferType()(mod)
+
+ in_data = np.random.randint(-120, high=120, size=ifm_shape, dtype="int8")
+ requant_data = (ifm_scale * (in_data - ifm_zp)) / ofm_scale + ofm_zp
+ out_data = np.round(np.clip(requant_data, -128, 127)).astype("int8")
+
+ compiled_model = infra.build_source(
+ mod, {"ifm": in_data}, [out_data], accel_type, output_tolerance=1
+ )
+
+ imported_modules = compiled_model[0].executor_factory.lib.imported_modules
+ assert len(imported_modules) == 2
+ ethosu_module = imported_modules[0]
+
+ # Verify generated C source
+ get_cs = tvm._ffi.get_global_func("runtime.module.ethosu.getcs")
+ cmms = get_cs(ethosu_module)
+ cmms = bytes.fromhex(cmms)
+
+ infra.print_payload(cmms)
+ infra.verify_source(compiled_model, accel_type)
+
+
+@pytest.mark.parametrize("accel_type", ACCEL_TYPES)
+@pytest.mark.parametrize(
+ "ifm_shape, new_shape",
+ [
+ ((1, 4, 1, 2), (1, 1, 1, 8)),
+ (
+ (
+ 5,
+ 1,
+ 20,
+ ),
+ (1, 5, 1, 20),
+ ),
+ ((12, 20), (1, 6, 4, 10)),
+ ((12, 20), (6, 4, 10)),
+ ((20,), (4, 5)),
+ ],
+)
+def test_relay_reshape_codegen(ifm_shape, new_shape, accel_type):
+ # Create a "partitioned" Relay graph
Review comment:
Done
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ekalda commented on a change in pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
Posted by GitBox <gi...@apache.org>.
ekalda commented on a change in pull request #9457:
URL: https://github.com/apache/tvm/pull/9457#discussion_r750115026
##########
File path: python/tvm/relay/backend/contrib/ethosu/legalize.py
##########
@@ -123,6 +123,108 @@ def __call__(self, *args, **kwargs):
pass
+class StridedSliceRewriter(DFPatternCallback):
+ """This pass brings the strided slice out of the partitioned function"""
+
+ def __init__(self):
+ super().__init__(require_type=True, rewrite_once=True)
+ self.pattern = (wildcard().has_attr({"Composite": "ethosu.strided_slice"}))(wildcard())
+
+ def callback(
+ self, pre: tvm.relay.Expr, post: tvm.relay.Expr, node_map: tvm.ir.container.Map
+ ) -> tvm.relay.Expr:
+ input = post.args[0]
+ attrs = post.op.body.attrs
+ begin = attrs.begin
+ end = attrs.end
+ strides = attrs.strides
+ axes = attrs.axes
+ slice_mode = attrs.slice_mode
+ strided_slice = relay.op.strided_slice(
+ input, begin, end, strides=strides, axes=axes, slice_mode=slice_mode
+ )
+ return strided_slice
+
+
+@ir.transform.module_pass(opt_level=1)
+class LegalizeStridedSlice:
+ """This is the pass that wraps StridedSliceRewriter"""
+
+ def transform_module(
+ self, mod: tvm.ir.IRModule, ctx: tvm.ir.transform.PassContext
+ ) -> tvm.ir.IRModule:
+ for global_var, func in mod.functions.items():
+ func = rewrite(StridedSliceRewriter(), func)
+ mod.update_func(global_var, func)
+ return mod
+
+ def __call__(self, *args, **kwargs):
+ pass
+
+
+class ReshapeRewriter(DFPatternCallback):
+ """This pass brings the reshape out of the partitioned function"""
+
+ def __init__(self):
+ super().__init__(require_type=True, rewrite_once=True)
+ self.pattern = (wildcard().has_attr({"Composite": "ethosu.reshape"}))(wildcard())
+
+ def callback(
+ self, pre: tvm.relay.Expr, post: tvm.relay.Expr, node_map: tvm.ir.container.Map
+ ) -> tvm.relay.Expr:
+ reshape_input = post.args[0]
+ new_shape = post.op.body.attrs.newshape
+ reshape = relay.op.reshape(reshape_input, newshape=new_shape)
+ return reshape
+
+
+@ir.transform.module_pass(opt_level=1)
+class LegalizeReshape:
+ """This is the pass that wraps ReshapeRewriter"""
+
+ def transform_module(
+ self, mod: tvm.ir.IRModule, ctx: tvm.ir.transform.PassContext
+ ) -> tvm.ir.IRModule:
+ for global_var, func in mod.functions.items():
+ func = rewrite(ReshapeRewriter(), func)
+ mod.update_func(global_var, func)
+ return mod
+
+ def __call__(self, *args, **kwargs):
+ pass
+
+
+class NoOpRewriter(DFPatternCallback):
+ """This pass adds and idenity operator to reshape and strided slice to avoid a no op without a consumer"""
Review comment:
Done
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ekalda commented on pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
Posted by GitBox <gi...@apache.org>.
ekalda commented on pull request #9457:
URL: https://github.com/apache/tvm/pull/9457#issuecomment-961802423
@manupa-arm @lhutton1 @dchauhan-arm @mbaret @NicolaLancellotti @Mousius
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ekalda commented on pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
Posted by GitBox <gi...@apache.org>.
ekalda commented on pull request #9457:
URL: https://github.com/apache/tvm/pull/9457#issuecomment-961802423
@manupa-arm @lhutton1 @dchauhan-arm @mbaret @NicolaLancellotti @Mousius
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] mbaret commented on a change in pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
Posted by GitBox <gi...@apache.org>.
mbaret commented on a change in pull request #9457:
URL: https://github.com/apache/tvm/pull/9457#discussion_r743789916
##########
File path: python/tvm/relay/backend/contrib/ethosu/op/identity.py
##########
@@ -0,0 +1,98 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+# pylint: disable=unused-argument
+"""Relay identity operator"""
+
+import tvm
+from tvm.relay.op import _make
+from tvm.topi.generic import schedule_injective
+from tvm.relay.op.op import OpStrategy
+from tvm.relay.op import strategy as _strategy
+
+from ..te import identity_compute
+
+
+@tvm.ir.register_op_attr("contrib.ethosu.identity", "FTVMCompute")
+def create_ethosu_identity_compute(attrs, args, out_type):
+ """Create an ethosu_identity compute op."""
+ ifm = args[0]
+ lut = args[1]
+ ifm_scale = attrs.ifm_scale
+ ifm_zero_point = attrs.ifm_zero_point
+ ofm_scale = attrs.ofm_scale
+ ofm_zero_point = attrs.ofm_zero_point
+ activation = attrs.activation
+ op = identity_compute(
+ ifm, lut, ifm_scale, ifm_zero_point, ofm_scale, ofm_zero_point, activation
+ )
+ return [op]
+
+
+@tvm.ir.register_op_attr("contrib.ethosu.identity", "FTVMStrategy")
+def pooling_strategy_ethosu(attrs, inputs, out_type, target):
Review comment:
identity_strategy_ethosu
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir/identity.py
##########
@@ -0,0 +1,87 @@
+# 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
+"""Extract information from the identity operator in TIR."""
+from typing import Dict, Tuple
+import tvm
+from .dma import get_read_params, get_write_params
+from .spec import SerialKernel, SerialActivation, SerialPooling, SerialPadding
+from .utils import get_op_attrs
+
+
+def get_identity_params(
+ stmt: tvm.tir.AttrStmt,
+ producers: Dict[tvm.tir.Var, tvm.tir.AttrStmt],
+ consumers: Dict[tvm.tir.Var, tvm.tir.AttrStmt],
+) -> Tuple[SerialPooling, tvm.tir.Var, tvm.tir.Var]:
+ """Get the parameters necessary to construct a call_extern for a pooling.
+
+ Parameters
+ ----------
+ stmt : tvm.tir.AttrStmt
+ The outermost attribute statement of a convolution loop nest.
+ producers : Dict[tvm.tir.Var, tvm.tir.AttrStmt]
+ A dictionary to associate pointers with the loop nest
+ that produces their values.
+ consumers : Dict[tvm.tir.Var, tvm.tir.AttrStmt]
+ A dictionary to associate pointers with the loop nest
+ that consumes their values.
+
+ Returns
+ -------
+ SerialPooling
+ The parameters needed to construct a 2D pooling.
+ output_pointer : tvm.tir.Var
+ The output pointer of the pooling operation.
+ replace_pointer : tvm.tir.Var
+ The output pointer of the DMA write operation, which is to replace
+ the pooling output pointer.
+
+ """
+ attrs, _ = get_op_attrs(stmt)
+ # Find the inner loop
+ while hasattr(stmt, "body"):
+ stmt = stmt.body
+
+ input_pointer = stmt.value.buffer_var
+ output_pointer = stmt.buffer_var
+
+ read = producers[input_pointer]
+ write = consumers[output_pointer]
+
+ serial_ifm, _, _ = get_read_params(read)
+ serial_ofm, _, write_output_pointer = get_write_params(write)
+
+ replace_pointer = write_output_pointer
+
+ # TODO (maybe): Support stand alone RELU through clamping in identity
Review comment:
Clarify this
##########
File path: tests/python/contrib/test_ethosu/test_type_inference.py
##########
@@ -171,5 +180,26 @@ def test_ethosu_pooling_invalid_dtype():
run_opt_pass(func, relay.transform.InferType())
+def test_ethosu_identity_invalid_shape():
+ invalid_shape = [1, 2, 3, 4, 5]
+ dtype = "int8"
+ ifm = relay.var("ifm", shape=invalid_shape, dtype=dtype)
+
+ identity = make_ethosu_identity(ifm)
+ func = relay.Function([ifm], identity)
+ with pytest.raises(TVMError):
+ run_opt_pass(func, relay.transform.InferType())
+
+
+def test_ethosu_invalid_invalid_dtype():
Review comment:
test_ethosu_identity_invalid_dtype
##########
File path: python/tvm/relay/backend/contrib/ethosu/legalize.py
##########
@@ -423,11 +526,15 @@ class LegalizeEthosU:
def transform_module(
self, mod: tvm.ir.IRModule, ctx: tvm.ir.transform.PassContext
) -> tvm.ir.IRModule:
+ """Legalize the oerators that can be offloaded to the NPU"""
Review comment:
operators
##########
File path: tests/python/contrib/test_ethosu/test_type_inference.py
##########
@@ -171,5 +180,26 @@ def test_ethosu_pooling_invalid_dtype():
run_opt_pass(func, relay.transform.InferType())
+def test_ethosu_identity_invalid_shape():
Review comment:
I think we need to test the valid cases as well.
##########
File path: python/tvm/relay/op/contrib/ethosu.py
##########
@@ -481,6 +493,8 @@ def pattern_table() -> List[Tuple[str, tvm.relay.dataflow_pattern.DFPattern, Cal
qnn_avgpool2d_pattern(),
lambda pat: AvgPool2DParams(pat).is_valid(),
),
+ ("ethosu.strided_slice", strided_slice_pattern(), lambda pat: True),
Review comment:
Let's put the proper restrictions here so we don't accidentally offload, for example, floating point tensors.
##########
File path: src/relay/op/contrib/ethosu/identity.cc
##########
@@ -0,0 +1,128 @@
+/*
+ * 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/op/contrib/ethosu/identity.cc
+ * \brief Property def of the Arm Ethos-U identity op.
+ */
+#include <tvm/relay/op.h>
+
+#include "common.h"
+
+namespace tvm {
+namespace relay {
+namespace op {
+namespace contrib {
+namespace ethosu {
+
+/*! \brief Attributes used by the Ethos(TM)-U NPU identity operator */
+struct EthosuIdentityAttrs : public tvm::AttrsNode<EthosuIdentityAttrs> {
+ double ifm_scale;
+ int ifm_zero_point;
+ double ofm_scale;
+ int ofm_zero_point;
+ String activation;
+
+ TVM_DECLARE_ATTRS(EthosuIdentityAttrs, "relay.attrs.EthosuIdentityAttrs") {
+ TVM_ATTR_FIELD(ifm_scale).describe("The quantization scale for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ifm_zero_point)
+ .describe("The quantization zero point for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_scale).describe("The quantization scale for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_zero_point)
+ .describe("The quantization zero point for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(activation)
+ .describe(
+ "The activation function to use. "
+ "'NONE' - no activation function. "
+ "'TANH' - tanh activation function. "
+ "'SIGMOID' - sigmoid activation function. "
+ "'LUT' - use a look-up table to perform the activation function.")
+ .set_default("NONE");
+ }
+};
+
+TVM_REGISTER_NODE_TYPE(EthosuIdentityAttrs);
+
+bool EthosuIdentityRel(const Array<Type>& types, int num_inputs, const Attrs& attrs,
+ const TypeReporter& reporter) {
+ int ifm_index = 0;
+ int result_index = 2;
+ ICHECK_EQ(types.size(), result_index + 1);
+
+ const auto* ifm = types[ifm_index].as<TensorTypeNode>();
+ if (ifm == nullptr) return false;
+
+ const auto* param = attrs.as<EthosuIdentityAttrs>();
+
+ ICHECK(param != nullptr) << "EthosuIdentityAttrs cannot be nullptr.";
+
+ if (ifm->dtype != DataType::UInt(8) && ifm->dtype != DataType::Int(8)) {
+ reporter->GetDiagCtx().EmitFatal(
+ Diagnostic::Error(reporter->GetSpan())
+ << "Invalid operator: Expected type(uint8) or type(int8) for ifm but was " << ifm->dtype);
+ return false;
+ }
+
+ if (ifm->shape.size() > 4) {
+ reporter->GetDiagCtx().EmitFatal(
+ Diagnostic::Error(reporter->GetSpan())
+ << "Invalid operator: Input Feature Map should be at most 4 dimensional, but was "
+ << ifm->shape);
+ return false;
+ }
+
+ // Assign ofm type
+ auto ofm_shape = ifm->shape;
+ reporter->Assign(types[result_index], TensorType(ofm_shape, ifm->dtype));
+ return true;
+}
+
+Expr MakeEthosuIdentity(Expr ifm, Expr lut, double ifm_scale, int ifm_zero_point, double ofm_scale,
+ int ofm_zero_point, String activation) {
+ auto attrs = make_object<EthosuIdentityAttrs>();
+ attrs->ifm_scale = ifm_scale;
+ attrs->ifm_zero_point = ifm_zero_point;
+ attrs->ofm_scale = ofm_scale;
+ attrs->ofm_zero_point = ofm_zero_point;
+ attrs->activation = std::move(activation);
+ static const Op& op = Op::Get("contrib.ethosu.identity");
+ return Call(op, {ifm, lut}, Attrs(attrs), {});
+}
+
+TVM_REGISTER_GLOBAL("relay.op._make.ethosu_identity").set_body_typed(MakeEthosuIdentity);
+
+RELAY_REGISTER_OP("contrib.ethosu.identity")
+ .describe(R"code(Identity operator for Ethos-U NPUs.
+
+This Relay operator performs the identity operation on Ethos(TM)-U NPU with a capability
+to requantize the data. It accepts input with any shape that is less or equal to 4.
Review comment:
Maybe 'It accepts input tensors of 4 dimensions or less.'?
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir/identity.py
##########
@@ -0,0 +1,87 @@
+# 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
+"""Extract information from the identity operator in TIR."""
+from typing import Dict, Tuple
+import tvm
+from .dma import get_read_params, get_write_params
+from .spec import SerialKernel, SerialActivation, SerialPooling, SerialPadding
+from .utils import get_op_attrs
+
+
+def get_identity_params(
+ stmt: tvm.tir.AttrStmt,
+ producers: Dict[tvm.tir.Var, tvm.tir.AttrStmt],
+ consumers: Dict[tvm.tir.Var, tvm.tir.AttrStmt],
+) -> Tuple[SerialPooling, tvm.tir.Var, tvm.tir.Var]:
+ """Get the parameters necessary to construct a call_extern for a pooling.
+
+ Parameters
+ ----------
+ stmt : tvm.tir.AttrStmt
+ The outermost attribute statement of a convolution loop nest.
Review comment:
identity pooling loop nest
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir/identity.py
##########
@@ -0,0 +1,87 @@
+# 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
+"""Extract information from the identity operator in TIR."""
+from typing import Dict, Tuple
+import tvm
+from .dma import get_read_params, get_write_params
+from .spec import SerialKernel, SerialActivation, SerialPooling, SerialPadding
+from .utils import get_op_attrs
+
+
+def get_identity_params(
+ stmt: tvm.tir.AttrStmt,
+ producers: Dict[tvm.tir.Var, tvm.tir.AttrStmt],
+ consumers: Dict[tvm.tir.Var, tvm.tir.AttrStmt],
+) -> Tuple[SerialPooling, tvm.tir.Var, tvm.tir.Var]:
+ """Get the parameters necessary to construct a call_extern for a pooling.
Review comment:
for an identity pooling
##########
File path: tests/python/contrib/test_ethosu/test_codegen.py
##########
@@ -343,5 +343,178 @@ def representative_dataset():
infra.verify_source(compiled_models, accel_type)
+@pytest.mark.parametrize("accel_type", ACCEL_TYPES)
+@pytest.mark.parametrize("ifm_shape", [(3, 2), (1, 15, 11, 7), (3, 1, 12), (400,)])
+@pytest.mark.parametrize("ifm_scale, ifm_zp, ofm_scale, ofm_zp", [(1, 0, 1, 0), (0.015, 3, 0.2, 5)])
+def test_ethosu_identity_codegen(ifm_shape, ifm_scale, ifm_zp, ofm_scale, ofm_zp, accel_type):
+ # Create a "partitioned" Relay function
+ ifm = relay.var("ifm", shape=ifm_shape, dtype="int8")
+ ifm0 = relay.var("ifm0", shape=ifm_shape, dtype="int8")
+ identity = infra.make_ethosu_identity(
+ ifm0, ifm_scale=ifm_scale, ifm_zero_point=ifm_zp, ofm_scale=ofm_scale, ofm_zero_point=ofm_zp
+ )
+ glb_ethosu = relay.GlobalVar("tvmgen_default_ethosu_main_0")
+
+ func = (
+ relay.Function([ifm0], identity)
+ .with_attr("Inline", 1)
+ .with_attr("Compiler", "ethosu")
+ .with_attr("global_symbol", "tvmgen_default_ethosu_main_0")
+ .with_attr("Primitive", 1)
+ )
+ mod = tvm.IRModule()
+ mod[glb_ethosu] = func
+ mod = relay.transform.InferType()(mod)
+
+ call = relay.Call(glb_ethosu, [ifm])
+ mod["main"] = relay.Function([ifm], call)
+ mod = relay.transform.InferType()(mod)
+
+ in_data = np.random.randint(-120, high=120, size=ifm_shape, dtype="int8")
+ requant_data = (ifm_scale * (in_data - ifm_zp)) / ofm_scale + ofm_zp
+ out_data = np.round(np.clip(requant_data, -128, 127)).astype("int8")
+
+ compiled_model = infra.build_source(
+ mod, {"ifm": in_data}, [out_data], accel_type, output_tolerance=1
+ )
+
+ imported_modules = compiled_model[0].executor_factory.lib.imported_modules
+ assert len(imported_modules) == 2
+ ethosu_module = imported_modules[0]
+
+ # Verify generated C source
+ get_cs = tvm._ffi.get_global_func("runtime.module.ethosu.getcs")
+ cmms = get_cs(ethosu_module)
+ cmms = bytes.fromhex(cmms)
+
+ infra.print_payload(cmms)
+ infra.verify_source(compiled_model, accel_type)
+
+
+@pytest.mark.parametrize("accel_type", ACCEL_TYPES)
+@pytest.mark.parametrize(
+ "ifm_shape, new_shape",
+ [
+ ((1, 4, 1, 2), (1, 1, 1, 8)),
+ (
+ (
+ 5,
+ 1,
+ 20,
+ ),
+ (1, 5, 1, 20),
+ ),
+ ((12, 20), (1, 6, 4, 10)),
+ ((12, 20), (6, 4, 10)),
+ ((20,), (4, 5)),
+ ],
+)
+def test_relay_reshape_codegen(ifm_shape, new_shape, accel_type):
+ # Create a "partitioned" Relay graph
Review comment:
It seems like some of this could be refactored out into a common function and reused between tests.
##########
File path: tests/python/contrib/test_ethosu/test_codegen.py
##########
@@ -343,5 +343,178 @@ def representative_dataset():
infra.verify_source(compiled_models, accel_type)
+@pytest.mark.parametrize("accel_type", ACCEL_TYPES)
+@pytest.mark.parametrize("ifm_shape", [(3, 2), (1, 15, 11, 7), (3, 1, 12), (400,)])
+@pytest.mark.parametrize("ifm_scale, ifm_zp, ofm_scale, ofm_zp", [(1, 0, 1, 0), (0.015, 3, 0.2, 5)])
+def test_ethosu_identity_codegen(ifm_shape, ifm_scale, ifm_zp, ofm_scale, ofm_zp, accel_type):
+ # Create a "partitioned" Relay function
+ ifm = relay.var("ifm", shape=ifm_shape, dtype="int8")
+ ifm0 = relay.var("ifm0", shape=ifm_shape, dtype="int8")
+ identity = infra.make_ethosu_identity(
+ ifm0, ifm_scale=ifm_scale, ifm_zero_point=ifm_zp, ofm_scale=ofm_scale, ofm_zero_point=ofm_zp
+ )
+ glb_ethosu = relay.GlobalVar("tvmgen_default_ethosu_main_0")
+
+ func = (
+ relay.Function([ifm0], identity)
+ .with_attr("Inline", 1)
+ .with_attr("Compiler", "ethosu")
+ .with_attr("global_symbol", "tvmgen_default_ethosu_main_0")
+ .with_attr("Primitive", 1)
+ )
+ mod = tvm.IRModule()
+ mod[glb_ethosu] = func
+ mod = relay.transform.InferType()(mod)
+
+ call = relay.Call(glb_ethosu, [ifm])
+ mod["main"] = relay.Function([ifm], call)
+ mod = relay.transform.InferType()(mod)
+
+ in_data = np.random.randint(-120, high=120, size=ifm_shape, dtype="int8")
+ requant_data = (ifm_scale * (in_data - ifm_zp)) / ofm_scale + ofm_zp
+ out_data = np.round(np.clip(requant_data, -128, 127)).astype("int8")
+
+ compiled_model = infra.build_source(
+ mod, {"ifm": in_data}, [out_data], accel_type, output_tolerance=1
+ )
+
+ imported_modules = compiled_model[0].executor_factory.lib.imported_modules
+ assert len(imported_modules) == 2
+ ethosu_module = imported_modules[0]
+
+ # Verify generated C source
+ get_cs = tvm._ffi.get_global_func("runtime.module.ethosu.getcs")
+ cmms = get_cs(ethosu_module)
+ cmms = bytes.fromhex(cmms)
+
+ infra.print_payload(cmms)
+ infra.verify_source(compiled_model, accel_type)
+
+
+@pytest.mark.parametrize("accel_type", ACCEL_TYPES)
+@pytest.mark.parametrize(
Review comment:
We should either test the 'special' indices for reshape (-1 and -2), or we should explicitly disallow these.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ekalda commented on pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
Posted by GitBox <gi...@apache.org>.
ekalda commented on pull request #9457:
URL: https://github.com/apache/tvm/pull/9457#issuecomment-961802423
@manupa-arm @lhutton1 @dchauhan-arm @mbaret @NicolaLancellotti @Mousius
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] ekalda commented on a change in pull request #9457: Add the Arm(R) Ethos(TM)-U NPU identity operator
Posted by GitBox <gi...@apache.org>.
ekalda commented on a change in pull request #9457:
URL: https://github.com/apache/tvm/pull/9457#discussion_r750120964
##########
File path: src/relay/op/contrib/ethosu/identity.cc
##########
@@ -0,0 +1,128 @@
+/*
+ * 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/op/contrib/ethosu/identity.cc
+ * \brief Property def of the Arm Ethos-U identity op.
+ */
+#include <tvm/relay/op.h>
+
+#include "common.h"
+
+namespace tvm {
+namespace relay {
+namespace op {
+namespace contrib {
+namespace ethosu {
+
+/*! \brief Attributes used by the Ethos(TM)-U NPU identity operator */
+struct EthosuIdentityAttrs : public tvm::AttrsNode<EthosuIdentityAttrs> {
+ double ifm_scale;
+ int ifm_zero_point;
+ double ofm_scale;
+ int ofm_zero_point;
+ String activation;
+
+ TVM_DECLARE_ATTRS(EthosuIdentityAttrs, "relay.attrs.EthosuIdentityAttrs") {
+ TVM_ATTR_FIELD(ifm_scale).describe("The quantization scale for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ifm_zero_point)
+ .describe("The quantization zero point for the Input Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_scale).describe("The quantization scale for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(ofm_zero_point)
+ .describe("The quantization zero point for the Output Feature Map tensor.");
+ TVM_ATTR_FIELD(activation)
+ .describe(
+ "The activation function to use. "
+ "'NONE' - no activation function. "
+ "'TANH' - tanh activation function. "
+ "'SIGMOID' - sigmoid activation function. "
+ "'LUT' - use a look-up table to perform the activation function.")
+ .set_default("NONE");
+ }
+};
+
+TVM_REGISTER_NODE_TYPE(EthosuIdentityAttrs);
+
+bool EthosuIdentityRel(const Array<Type>& types, int num_inputs, const Attrs& attrs,
+ const TypeReporter& reporter) {
+ int ifm_index = 0;
+ int result_index = 2;
Review comment:
Done
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org