Arm(R) Ethos(TM)-U NPU Depthwise2d operator support by ekalda · Pull Request #9209 · apache/tvm

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

manupak

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

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



+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

+ 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

+ 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

+ 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

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


nit : type annotations

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

+ 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

 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

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

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

 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

+ 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

+ 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

+ # 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 · 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

manupak

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

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



+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

+ 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

+ 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

+ 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

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


nit : type annotations

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

+ 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

 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

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

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

 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

+ 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

+ 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

+ # 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 · 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

manupak

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

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



+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

+ 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

+ 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

+ 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

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


nit : type annotations

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

+ 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

 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

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

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

 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

+ 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

+ 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

+ # 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 · 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

manupak

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

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



+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

+ 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

+ 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

+ 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

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


nit : type annotations

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

+ 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

 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

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

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

 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

+ 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

+ 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

+ # 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 · 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

manupak

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

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



+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

+ 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

+ 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

+ 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

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


nit : type annotations

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

+ 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

 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

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

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

 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

+ 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

+ 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

+ # 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 · 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

manupak

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

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



+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

+ 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

+ 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

+ 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

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


nit : type annotations

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

+ 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

 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

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

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

 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

+ 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

+ 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

+ # 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 · 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

manupak

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

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



+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

+ 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

+ 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

+ 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

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


nit : type annotations

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

+ 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

 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

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

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

 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

+ 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

+ 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

+ # 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 · 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

manupak

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

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



+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

+ 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

+ 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

+ 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

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


nit : type annotations

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

+ 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

 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

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

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

 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

+ 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

+ 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

+ # 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 MarisaKirisame, ZihengJiang, anijain2305, areusch, comaniac, jroesch, junrushao, jwfromm, mbrookhart, merrymercy, slyubomirsky, tqchen, vinx13, wweic, yzhliu and zhiics as code owners October 6, 2021 14:42

mbaret reviewed Oct 6, 2021
View reviewed changes

Comment threadtests/python/contrib/test_ethosu/test_legalize.py Outdated

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 1e7e9caCompare 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 f43e088Compare October 11, 2021 09:08

mbaret approved these changes Oct 11, 2021
View reviewed changes

mbaret merged commit 8ba0451 into apache:mainOct 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

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

mbaret reviewed Oct 6, 2021
View reviewed changes

Comment threadtests/python/contrib/test_ethosu/test_legalize.py Outdated

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 1e7e9caCompare 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 f43e088Compare October 11, 2021 09:08

mbaret approved these changes Oct 11, 2021
View reviewed changes

mbaret merged commit 8ba0451 into apache:mainOct 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

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

mbaret reviewed Oct 6, 2021
View reviewed changes

Comment threadtests/python/contrib/test_ethosu/test_legalize.py Outdated

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 1e7e9caCompare 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 f43e088Compare October 11, 2021 09:08

mbaret approved these changes Oct 11, 2021
View reviewed changes

mbaret merged commit 8ba0451 into apache:mainOct 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

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

mbaret reviewed Oct 6, 2021
View reviewed changes

Comment threadtests/python/contrib/test_ethosu/test_legalize.py Outdated

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 1e7e9caCompare 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 f43e088Compare October 11, 2021 09:08

mbaret approved these changes Oct 11, 2021
View reviewed changes

mbaret merged commit 8ba0451 into apache:mainOct 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

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

mbaret reviewed Oct 6, 2021
View reviewed changes

Comment threadtests/python/contrib/test_ethosu/test_legalize.py Outdated

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 1e7e9caCompare 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 f43e088Compare October 11, 2021 09:08

mbaret approved these changes Oct 11, 2021
View reviewed changes

mbaret merged commit 8ba0451 into apache:mainOct 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

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

mbaret reviewed Oct 6, 2021
View reviewed changes

Comment threadtests/python/contrib/test_ethosu/test_legalize.py Outdated

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 1e7e9caCompare 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 f43e088Compare October 11, 2021 09:08

mbaret approved these changes Oct 11, 2021
View reviewed changes

mbaret merged commit 8ba0451 into apache:mainOct 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

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

mbaret reviewed Oct 6, 2021
View reviewed changes

Comment threadtests/python/contrib/test_ethosu/test_legalize.py Outdated

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 1e7e9caCompare 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 f43e088Compare October 11, 2021 09:08

mbaret approved these changes Oct 11, 2021
View reviewed changes

mbaret merged commit 8ba0451 into apache:mainOct 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

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

mbaret reviewed Oct 6, 2021
View reviewed changes

Comment threadtests/python/contrib/test_ethosu/test_legalize.py Outdated

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 1e7e9caCompare 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 f43e088Compare October 11, 2021 09:08

mbaret approved these changes Oct 11, 2021
View reviewed changes

mbaret merged commit 8ba0451 into apache:mainOct 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



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

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

	def__init__(self, func_body):
	def__init__(self, func_body: tvm.relay.expr.Call):

Uh oh!

Conversation

ekalda commented Oct 6, 2021

Uh oh!

ekalda commented Oct 6, 2021

Uh oh!

manupak commented Oct 6, 2021

Uh oh!

Uh oh!

manupak left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

manupak commented Oct 7, 2021

Uh oh!

lhutton1 left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

manupak left a comment

Choose a reason for hiding this comment

Uh oh!

manupakOct 8, 2021• edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

lhutton1 left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

manupak left a comment • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

manupakOct 8, 2021•
edited
Loading

manupak left a comment •
edited
Loading