Arm(R) Ethos(TM)-U NPU Depthwise2d operator support #9209

ekalda · 2021-10-06T14:42:20Z

This commit adds support for Depthwise2d primitive operator throughout
the TVM stack including Relay legalization pass, operator definition,
TE, TIR passes and translation into the command stream.

ekalda · 2021-10-06T14:43:43Z

@manupa-arm @mbaret

manupak · 2021-10-06T14:47:33Z

Thanks @ekalda!

Just a high-level comment, I think we should stick to DepthwiseConv2D (as opposed to Depthwise2D). I ll have a look,

Also cc : @NicolaLancellotti @lhutton1

tests/python/contrib/test_ethosu/test_legalize.py

manupak

@ekalda I have left some minor comments, otherwise the implementation generally looks good.

manupak · 2021-10-07T05:48:38Z

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

@@ -208,6 +208,96 @@ def __call__(self, *args, **kwargs):
        pass


+class EthosuDepthwise2DRewriter(DFPatternCallback):
+    """Convert ethosu.qnn_depthwise2d composite functions to ethosu_depthwise2d operators"""


Let us stick to depthwiseconv2d/DepthwiseConv2D and also in the following mentions to it.

Done (I used depthwise_conv2d since it is more readable, but I can change it to depthwiseconv2d if you'd prefer that)

manupak · 2021-10-07T05:49:04Z

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

+    def __init__(self):
+        super().__init__(require_type=True)
+        self.pattern = (
+            wildcard().has_attr({"Composite": ethosu_patterns.QnnDepthwise2DParams.composite_name})


QnnDepthwiseConv2DParams

manupak · 2021-10-07T05:51:15Z

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

+    ofm_zero_point: int,
+    kernel_shape: Tuple[int, int],
+    ofm_channels: int,
+    strides: Tuple[int, int] = (1, 1),


nit : We can use Optional[Tuple[int, int]]

I left it as is since Optional is used when the variable can take a value None

manupak · 2021-10-07T05:52:13Z

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

+        The OFM tensor.
+
+    """
+    assert ifm.shape[0] == 1


It is better to give a message when this fails as to why it was assumed to be 1.

manupak · 2021-10-07T05:53:11Z

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

+)
+
+
+def get_depthwise2d_params(stmt, producers, consumers):


nit : type annotations

tests/python/contrib/test_ethosu/test_legalize.py

manupak · 2021-10-07T06:05:53Z

also cc : @dchauhan-arm

lhutton1

LGTM, just a couple of questions/minor things

lhutton1 · 2021-10-06T21:28:38Z

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

+        if str(params.ofm.layout) not in channels_map.keys():
+            raise UnsupportedLayout(str(params.ofm.layout))
+        kernel_shape_map = {
+            "HWOI": params.weights.shape[0:2],


Is it worth supporting OHWI weights here?

IIRC, in the Relay that corresponds to depthwise conv2d operator from TFLite, the weights are always in HWOI, that's why other formats are not handled here.

lhutton1 · 2021-10-06T21:58:23Z

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

@@ -223,12 +223,13 @@ def __init__(self, func_body: tvm.relay.Function):
        self.strides = qnn_conv2d.attrs.strides
        self.dilation = qnn_conv2d.attrs.dilation
        self.activation = activation
+        self.channels = qnn_conv2d.attrs.channels


Better to access attrs once here i.e,

attrs = qnn_conv2d.attrs self.padding = attrs.padding ... self.channels = attrs.channels

lhutton1 · 2021-10-06T22:05:43Z

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

+    ).has_attr({"kernel_layout": "HWOI"})
+    bias_add = is_op("nn.bias_add")(qnn_conv2d, is_constant())
+    req = is_op("qnn.requantize")(
+        qnn_conv2d | bias_add, is_constant(), is_constant(), is_constant(), is_constant()


Remove optional bias here? Then we can follow up with separate PR for conv2d

Good point, done

NicolaLancellotti · 2021-10-06T16:29:33Z

tests/python/contrib/test_ethosu/test_legalize.py

+
+    depthwise_pattern_table = [
+        (
+            "ethosu.depthwise2d",


Suggested change

"ethosu.depthwise2d",

"ethosu.QnnDepthwise2DParams.composite_name",

manupak

Just one comment otherwise LGTM modulo others' comments.

manupak · 2021-10-08T09:14:16Z

tests/python/driver/tvmc/test_compiler.py

@@ -471,7 +471,7 @@ def test_compile_tflite_module_with_external_codegen_ethosu(
                for name in mlf_package.getnames()
                if re.match(r"\./codegen/host/src/\D+\d+\.c", name)
            ]
-            assert len(c_source_files) == 17
+            assert len(c_source_files) == 4


We should have put a comment saying why there was 17 here originally. Sorry about that.
Would you be able to put a comments explaining why it is 4 now ?

It should along the lines of that we expect lesser subgraphs where it was just conv2D being offloaded and now we have depthwise_conv2d being offloaded as well from mobilenet. Therefore [conv2d-->dethpwise_conv2d-->conv2d-> ... ] get fused to a single primitive external function.

Added a comment, does it make sense?

lhutton1

LGTM

dchauhan-arm · 2021-10-08T10:18:27Z

tests/python/contrib/test_ethosu/test_codegen.py

+                if activation:
+                    op = tf.nn.relu(op)
+                return op
+


LGTM regarding tf.nn.depthwise_conv2d usage. Others' comments cover everything else.

NicolaLancellotti · 2021-10-08T10:47:13Z

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

+    stmt: tvm.tir.AttrStmt,
+    producers: Dict[tvm.tir.Var, tvm.tir.AttrStmt],
+    consumers: Dict[tvm.tir.Var, tvm.tir.AttrStmt],
+):


Suggested change

):

) -> Tuple[SerialPooling, tvm.tir.Var, tvm.tir.Var]:

NicolaLancellotti · 2021-10-08T10:57:55Z

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

+    # The hardware only supports padding upto the numbers as follows
+    padding_bounds = [31, 31, 32, 32]
+
+    def __init__(self, func_body):


Suggested change

def __init__(self, func_body):

def __init__(self, func_body: tvm.relay.expr.Call):

manupak

LGTM.

Other reviewers please approve explicitly if the discussions are resolved.
https://tvm.apache.org/docs/contribute/code_review.html#approve-and-request-changes-explicitly

This commit adds support for Depthwise2d primitive operator throughout the TVM stack including Relay legalization pass, operator definition, TE, TIR passes and translation into the command stream. Change-Id: If82b85f5d3b23cd214fe38babd724451bf95ef5b

And respond to other review comments. Change-Id: I58a9f28723750970d386b4d0ba62fa399c5c6181

Change-Id: Idf4c078bf65e7ed31fe82a92bf334295a82b6ead

Change-Id: Ic6c77af30a5b9cb68dcc0c173b95490965359481

Change-Id: I7318bd8cfa5985b33fc7d020cc19057cc9498197

mbaret

LGTM, let's get this in.

mbaret · 2021-10-11T15:38:04Z

This is now merged, thanks everyone!

* Arm(R) Ethos(TM)-U NPU Depthwise2d operator support This commit adds support for Depthwise2d primitive operator throughout the TVM stack including Relay legalization pass, operator definition, TE, TIR passes and translation into the command stream. Change-Id: If82b85f5d3b23cd214fe38babd724451bf95ef5b * Change depthwise2d to depthwise_conv2d And respond to other review comments. Change-Id: I58a9f28723750970d386b4d0ba62fa399c5c6181 * Make a line shorter and add a comment Change-Id: Idf4c078bf65e7ed31fe82a92bf334295a82b6ead * Change the order of imports Change-Id: Ic6c77af30a5b9cb68dcc0c173b95490965359481 * Whitespace change Change-Id: I7318bd8cfa5985b33fc7d020cc19057cc9498197

ekalda requested review from anijain2305, areusch, comaniac, jroesch, junrushao, jwfromm, MarisaKirisame, mbrookhart, merrymercy, slyubomirsky, tqchen, vinx13, wweic, yzhliu, zhiics and ZihengJiang as code owners October 6, 2021 14:42

mbaret reviewed Oct 6, 2021

View reviewed changes

tests/python/contrib/test_ethosu/test_legalize.py Outdated Show resolved Hide resolved

manupak requested changes Oct 7, 2021

View reviewed changes

lhutton1 reviewed Oct 7, 2021

View reviewed changes

NicolaLancellotti reviewed Oct 7, 2021

View reviewed changes

ekalda force-pushed the depthwise_upstream2 branch from 0812f86 to 1e7e9ca Compare October 8, 2021 08:57

manupak requested changes Oct 8, 2021

View reviewed changes

manupak reviewed Oct 8, 2021

View reviewed changes

lhutton1 approved these changes Oct 8, 2021

View reviewed changes

dchauhan-arm reviewed Oct 8, 2021

View reviewed changes

NicolaLancellotti reviewed Oct 8, 2021

View reviewed changes

manupak approved these changes Oct 8, 2021

View reviewed changes

manupak self-assigned this Oct 8, 2021

NicolaLancellotti approved these changes Oct 8, 2021

View reviewed changes

ekalda added 5 commits October 11, 2021 10:04

Change depthwise2d to depthwise_conv2d

1909515

And respond to other review comments. Change-Id: I58a9f28723750970d386b4d0ba62fa399c5c6181

Make a line shorter and add a comment

7c510fc

Change-Id: Idf4c078bf65e7ed31fe82a92bf334295a82b6ead

Change the order of imports

1d586dd

Change-Id: Ic6c77af30a5b9cb68dcc0c173b95490965359481

Whitespace change

f43e088

Change-Id: I7318bd8cfa5985b33fc7d020cc19057cc9498197

ekalda force-pushed the depthwise_upstream2 branch from 76ceb71 to f43e088 Compare October 11, 2021 09:08

mbaret approved these changes Oct 11, 2021

View reviewed changes

mbaret merged commit 8ba0451 into apache:main Oct 11, 2021

ekalda deleted the depthwise_upstream2 branch October 11, 2021 15:51

junrushao mentioned this pull request Nov 2, 2021

Apache TVM v0.8 Release Note Candidate #9416

Closed

	"ethosu.depthwise2d",
	"ethosu.QnnDepthwise2DParams.composite_name",

	def __init__(self, func_body):
	def __init__(self, func_body: tvm.relay.expr.Call):

Arm(R) Ethos(TM)-U NPU Depthwise2d operator support #9209

Arm(R) Ethos(TM)-U NPU Depthwise2d operator support #9209

Conversation

ekalda commented Oct 6, 2021

ekalda commented Oct 6, 2021

manupak commented Oct 6, 2021

manupak left a 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

manupak commented Oct 7, 2021

lhutton1 left a 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

manupak left a comment

Choose a reason for hiding this comment

manupak Oct 8, 2021 • edited Loading

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

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

manupak left a comment • edited Loading

Choose a reason for hiding this comment

mbaret left a comment

Choose a reason for hiding this comment

mbaret commented Oct 11, 2021

manupak Oct 8, 2021 •

edited

Loading

manupak left a comment •

edited

Loading