Add the Arm(R) Ethos(TM)-U NPU identity operator #9457

ekalda · 2021-11-05T11:00:09Z

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

ekalda · 2021-11-05T11:02:30Z

@manupa-arm @lhutton1 @dchauhan-arm @mbaret @NicolaLancellotti @Mousius

python/tvm/relay/backend/contrib/ethosu/legalize.py

python/tvm/relay/backend/contrib/ethosu/te/dma.py

python/tvm/relay/backend/contrib/ethosu/op/identity.py

mbaret · 2021-11-05T16:25:16Z

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"""


A different docstring appeared in the rebase, so I kept that that different docstring

mbaret · 2021-11-05T16:33:47Z

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),


Let's put the proper restrictions here so we don't accidentally offload, for example, floating point tensors.

python/tvm/relay/backend/contrib/ethosu/tir/identity.py

mbaret · 2021-11-05T16:56:57Z

python/tvm/relay/backend/contrib/ethosu/tir/identity.py

+    Parameters
+    ----------
+    stmt : tvm.tir.AttrStmt
+        The outermost attribute statement of a convolution loop nest.


identity pooling loop nest

mbaret · 2021-11-05T17:01:08Z

src/relay/op/contrib/ethosu/identity.cc

+    .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.


Maybe 'It accepts input tensors of 4 dimensions or less.'?

mbaret · 2021-11-05T17:04:39Z

tests/python/contrib/test_ethosu/test_codegen.py

+
+
+@pytest.mark.parametrize("accel_type", ACCEL_TYPES)
+@pytest.mark.parametrize(


We should either test the 'special' indices for reshape (-1 and -2), or we should explicitly disallow these.

mbaret · 2021-11-05T17:05:19Z

tests/python/contrib/test_ethosu/test_codegen.py

+    ],
+)
+def test_relay_reshape_codegen(ifm_shape, new_shape, accel_type):
+    # Create a "partitioned" Relay graph


It seems like some of this could be refactored out into a common function and reused between tests.

mbaret · 2021-11-05T17:15:45Z

tests/python/contrib/test_ethosu/test_type_inference.py

+        run_opt_pass(func, relay.transform.InferType())
+
+
+def test_ethosu_invalid_invalid_dtype():


test_ethosu_identity_invalid_dtype

mbaret · 2021-11-05T17:16:17Z

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():


I think we need to test the valid cases as well.

lhutton1

Thanks @ekalda, great work :) Just a few small things below..

python/tvm/relay/backend/contrib/ethosu/legalize.py

python/tvm/relay/backend/contrib/ethosu/te/identity.py

src/relay/op/contrib/ethosu/identity.cc

tests/python/contrib/test_ethosu/test_type_inference.py

* 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 Change-Id: I0adb5ca269f8529c79e0e7681ca4b5147d8f53c8

Change-Id: Icc9b6507f164681a5d6b1fcff2ae4a5051d44734

Change-Id: I63f30f84ad481789fc047ad8c2107f5313562f7f

ekalda · 2021-11-16T10:23:19Z

@mbaret @NicolaLancellotti @lhutton1 thanks for the reviews, much appreciated! :)

NicolaLancellotti

LGTM

manupak

Thanks @ekalda . LGTM!

@mbaret, @lhutton1 could you take an another look ?

lhutton1

LGTM, thanks @ekalda!

mbaret · 2021-11-16T17:25:08Z

Thanks everyone, this is now merged!

* Add the Arm(R) Ethos(TM)-U NPU identity operator * 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 Change-Id: I0adb5ca269f8529c79e0e7681ca4b5147d8f53c8 * Fix the pylint errors Change-Id: Icc9b6507f164681a5d6b1fcff2ae4a5051d44734 * Changes in response to review comments Change-Id: I63f30f84ad481789fc047ad8c2107f5313562f7f

ekalda requested review from anijain2305, areusch, comaniac, icemelon, jroesch, junrushao, jwfromm, MarisaKirisame, mbrookhart, merrymercy, slyubomirsky, tqchen, vinx13, wweic, yzhliu, zhiics and ZihengJiang as code owners November 5, 2021 11:00

NicolaLancellotti reviewed Nov 5, 2021

View reviewed changes

python/tvm/relay/backend/contrib/ethosu/legalize.py Outdated Show resolved Hide resolved

python/tvm/relay/backend/contrib/ethosu/te/dma.py Outdated Show resolved Hide resolved

python/tvm/relay/backend/contrib/ethosu/te/dma.py Outdated Show resolved Hide resolved

mbaret reviewed Nov 5, 2021

View reviewed changes

ekalda mentioned this pull request Nov 8, 2021

[Tracking Issue] Arm(R) Ethos(TM)-U NPU operator support #9410

Closed

15 tasks

lhutton1 reviewed Nov 8, 2021

View reviewed changes

ekalda added 3 commits November 15, 2021 12:06

Fix the pylint errors

9630151

Change-Id: Icc9b6507f164681a5d6b1fcff2ae4a5051d44734

Changes in response to review comments

ac87b0b

Change-Id: I63f30f84ad481789fc047ad8c2107f5313562f7f

ekalda force-pushed the identity_upstream branch from 75d1117 to ac87b0b Compare November 16, 2021 10:12

NicolaLancellotti approved these changes Nov 16, 2021

View reviewed changes

manupak approved these changes Nov 16, 2021

View reviewed changes

lhutton1 approved these changes Nov 16, 2021

View reviewed changes

mbaret merged commit c3a06ee into apache:main Nov 16, 2021

driazati mentioned this pull request Jul 14, 2022

TVM v0.9.0.rc0 Release Candidate Notes #12102

Closed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Add the Arm(R) Ethos(TM)-U NPU identity operator #9457

Add the Arm(R) Ethos(TM)-U NPU identity operator #9457

ekalda commented Nov 5, 2021

ekalda commented Nov 5, 2021

mbaret Nov 5, 2021

ekalda Nov 16, 2021

mbaret Nov 5, 2021

ekalda Nov 16, 2021

mbaret Nov 5, 2021

ekalda Nov 16, 2021

mbaret Nov 5, 2021

ekalda Nov 16, 2021

mbaret Nov 5, 2021

ekalda Nov 16, 2021

mbaret Nov 5, 2021

ekalda Nov 16, 2021

mbaret Nov 5, 2021

ekalda Nov 16, 2021

mbaret Nov 5, 2021

ekalda Nov 16, 2021

lhutton1 left a comment

ekalda commented Nov 16, 2021

NicolaLancellotti left a comment

manupak left a comment •

edited

Loading

lhutton1 left a comment

mbaret commented Nov 16, 2021



		@pytest.mark.parametrize("accel_type", ACCEL_TYPES)
		@pytest.mark.parametrize(

		run_opt_pass(func, relay.transform.InferType())


		def test_ethosu_invalid_invalid_dtype():

		@@ -171,5 +180,26 @@ def test_ethosu_pooling_invalid_dtype():
		run_opt_pass(func, relay.transform.InferType())


		def test_ethosu_identity_invalid_shape():

Add the Arm(R) Ethos(TM)-U NPU identity operator #9457

Add the Arm(R) Ethos(TM)-U NPU identity operator #9457

Conversation

ekalda commented Nov 5, 2021

ekalda commented Nov 5, 2021

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

lhutton1 left a comment

Choose a reason for hiding this comment

ekalda commented Nov 16, 2021

NicolaLancellotti left a comment

Choose a reason for hiding this comment

manupak left a comment • edited Loading

Choose a reason for hiding this comment

lhutton1 left a comment

Choose a reason for hiding this comment

mbaret commented Nov 16, 2021

manupak left a comment •

edited

Loading