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 2022/05/25 09:52:29 UTC
[GitHub] [tvm] NicolaLancellotti opened a new pull request, #11453: [microNPU] Add transform matrices and part matcher to identity op
NicolaLancellotti opened a new pull request, #11453:
URL: https://github.com/apache/tvm/pull/11453
This commit adds support for the identity operator in the cascader.
@manupa-arm @ekalda @lhutton1
--
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 diff in pull request #11453: [microNPU] Add transform matrices and part matcher to identity op
Posted by GitBox <gi...@apache.org>.
NicolaLancellotti commented on code in PR #11453:
URL: https://github.com/apache/tvm/pull/11453#discussion_r882847469
##########
python/tvm/relay/backend/contrib/ethosu/te/identity.py:
##########
@@ -76,7 +78,85 @@ def identity_compute(
name="ethosu_identity",
attrs=id_attrs,
)
+ length = len(ifm.shape)
+ ifm_matrix = np.identity(length + 1)
+ offset = np.zeros(length, dtype="int64")
+ ifm_propagator = Propagator(
+ ifm_matrix,
+ offset.tolist(),
+ )
+ propagator_attrs = {
+ "ifm_propagator": ifm_propagator,
+ }
+ return write_compute(identity, ofm_zero_point, ofm_scale, attrs=propagator_attrs)
+
+
+@register_matcher
+def match_ethosu_identity(output_tensor, device_config):
+ """Match a Tensor Expression corresponding to an NPU identity.
- dmaed_ofm = write_compute(identity, ofm_zero_point, ofm_scale)
+ If the Tensor Expression matches, an EthosuPart will be created that models the
+ matched Tensor Expression. Otherwise, None will be returned.
- return dmaed_ofm
+ Parameters
+ ----------
+ output_tensor : tvm.te.Tensor
+ The tensor to attempt to match with.
+ device_config : EthosuDeviceConfig
+ Target device configuration
+
+ Returns
+ -------
+ Union[None, EthosuPart]
+ The created EthosuPart if there was a match, otherwise None.
+ """
+ write = output_tensor
+ if write.op.name != "ethosu_write":
+ return None
+ identity = write.op.input_tensors[0]
+ if identity.op.name != "ethosu_identity":
+ return None
+ read = identity.op.input_tensors[0]
+ if read.op.name != "ethosu_read":
+ return None
+
+ input_tensors = [
+ read.op.input_tensors[0],
+ ]
+ subgraph = TESubgraph(input_tensors, output_tensor)
+ propagators = [
+ write.op.attrs["ifm_propagator"],
+ ]
+ ifm_dtype = input_tensors[0].dtype
+ ofm_dtype = output_tensor.dtype
+
+ input_tensors_shape = input_tensors[0].shape
+ ifm_channels = int(input_tensors_shape[3] if len(input_tensors_shape) > 3 else 1)
+ ofm_channels = ifm_channels
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 #11453: [microNPU] Add transform matrices and part matcher to identity op
Posted by GitBox <gi...@apache.org>.
ekalda commented on PR #11453:
URL: https://github.com/apache/tvm/pull/11453#issuecomment-1138553450
(I don't know why it displays each comment twice :sweat_smile: )
--
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 pull request #11453: [microNPU] Add transform matrices and part matcher to identity op
Posted by GitBox <gi...@apache.org>.
NicolaLancellotti commented on PR #11453:
URL: https://github.com/apache/tvm/pull/11453#issuecomment-1138748743
>
> (I don't know why it displays each comment twice )
No, problem. Thanks, @ekalda for the review.
--
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 diff in pull request #11453: [microNPU] Add transform matrices and part matcher to identity op
Posted by GitBox <gi...@apache.org>.
NicolaLancellotti commented on code in PR #11453:
URL: https://github.com/apache/tvm/pull/11453#discussion_r882846863
##########
python/tvm/contrib/ethosu/cascader/device_config.py:
##########
@@ -609,18 +624,19 @@ def _get_subkernel_propagator(
stride_w = int(op_attrs.get("stride_w", 1))
transform = ifm_propagator.transform
- if input_layout == "NHCWB16":
- transform[1][-1] = min(transform[1][-1], self._subkernel_limits[0] - stride_h)
- transform[3][-1] = min(transform[3][-1], self._subkernel_limits[1] - stride_w)
- else:
- transform[1][-1] = min(transform[1][-1], self._subkernel_limits[0] - stride_h)
- transform[2][-1] = min(transform[2][-1], self._subkernel_limits[1] - stride_w)
-
- if op_type in ("ethosu_pooling", "ethosu_depthwise_conv2d"):
- if output_layout == "NHCWB16" and input_layout == "NHWC":
- transform[3][-1] = depth
- elif output_layout == "NHCWB16" and input_layout == "NHCWB16":
- transform[2][-1] = 1 + ((depth - 1) // 16)
+ if op_type != "ethosu_identity":
Review Comment:
The block assumes that the transform matrix has at least three rows, but the one for identity can have fewer rows, so in that case, it crashes.
It is safe to hide that block because, for the identity, the transform matrix has only 0 and 1 as values, both less than the second arguments of the min function, so that block does not change the transform matrix.
--
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 merged pull request #11453: [microNPU] Add transform matrices and part matcher to identity op
Posted by GitBox <gi...@apache.org>.
lhutton1 merged PR #11453:
URL: https://github.com/apache/tvm/pull/11453
--
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 diff in pull request #11453: [microNPU] Add transform matrices and part matcher to identity op
Posted by GitBox <gi...@apache.org>.
NicolaLancellotti commented on code in PR #11453:
URL: https://github.com/apache/tvm/pull/11453#discussion_r885654969
##########
python/tvm/relay/backend/contrib/ethosu/te/identity.py:
##########
@@ -76,7 +78,85 @@ def identity_compute(
name="ethosu_identity",
attrs=id_attrs,
)
+ length = len(ifm.shape)
+ ifm_matrix = np.identity(length + 1)
+ offset = np.zeros(length, dtype="int64")
+ ifm_propagator = Propagator(
+ ifm_matrix,
+ offset.tolist(),
Review Comment:
You are welcome.
##########
python/tvm/relay/backend/contrib/ethosu/te/identity.py:
##########
@@ -76,7 +78,85 @@ def identity_compute(
name="ethosu_identity",
attrs=id_attrs,
)
+ length = len(ifm.shape)
+ ifm_matrix = np.identity(length + 1)
+ offset = np.zeros(length, dtype="int64")
+ ifm_propagator = Propagator(
+ ifm_matrix,
+ offset.tolist(),
+ )
+ propagator_attrs = {
+ "ifm_propagator": ifm_propagator,
+ }
+ return write_compute(identity, ofm_zero_point, ofm_scale, attrs=propagator_attrs)
+
+
+@register_matcher
+def match_ethosu_identity(output_tensor, device_config):
+ """Match a Tensor Expression corresponding to an NPU identity.
- dmaed_ofm = write_compute(identity, ofm_zero_point, ofm_scale)
+ If the Tensor Expression matches, an EthosuPart will be created that models the
+ matched Tensor Expression. Otherwise, None will be returned.
- return dmaed_ofm
+ Parameters
+ ----------
+ output_tensor : tvm.te.Tensor
+ The tensor to attempt to match with.
+ device_config : EthosuDeviceConfig
+ Target device configuration
+
+ Returns
+ -------
+ Union[None, EthosuPart]
+ The created EthosuPart if there was a match, otherwise None.
+ """
+ write = output_tensor
+ if write.op.name != "ethosu_write":
+ return None
+ identity = write.op.input_tensors[0]
+ if identity.op.name != "ethosu_identity":
+ return None
+ read = identity.op.input_tensors[0]
+ if read.op.name != "ethosu_read":
+ return None
+
+ input_tensors = [
+ read.op.input_tensors[0],
+ ]
+ subgraph = TESubgraph(input_tensors, output_tensor)
+ propagators = [
+ write.op.attrs["ifm_propagator"],
+ ]
+ ifm_dtype = input_tensors[0].dtype
+ ofm_dtype = output_tensor.dtype
+
+ input_tensors_shape = input_tensors[0].shape
+ length = len(input_tensors_shape)
+ channels = int(input_tensors_shape[length - 1]) if length >= 3 else 1
+
+ subkernels = len(device_config.get_kernel_steps(identity.op.name, 1, 1, ifm_dtype))
+
+ input_layout = output_layout = "NHWC"
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 diff in pull request #11453: [microNPU] Add transform matrices and part matcher to identity op
Posted by GitBox <gi...@apache.org>.
NicolaLancellotti commented on code in PR #11453:
URL: https://github.com/apache/tvm/pull/11453#discussion_r885646324
##########
python/tvm/relay/backend/contrib/ethosu/te/identity.py:
##########
@@ -76,7 +78,85 @@ def identity_compute(
name="ethosu_identity",
attrs=id_attrs,
)
+ length = len(ifm.shape)
+ ifm_matrix = np.identity(length + 1)
+ offset = np.zeros(length, dtype="int64")
+ ifm_propagator = Propagator(
+ ifm_matrix,
+ offset.tolist(),
Review Comment:
The offset must be a python list, instead, the numpy transform matrix is converted to a list by the Propagator constructor. https://github.com/apache/tvm/blob/efec735626cb2932df7771bdff60c2e9ff1b14da/python/tvm/contrib/ethosu/cascader/propagator.py#L30-L32
--
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] manupa-arm commented on pull request #11453: [microNPU] Add transform matrices and part matcher to identity op
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on PR #11453:
URL: https://github.com/apache/tvm/pull/11453#issuecomment-1140885316
Hi @NicolaLancellotti ,
It would be great to rebase this PR on top of https://github.com/apache/tvm/pull/11410/files and see if the unit tests that I have disabled in that PR works -- so we can have higher confidence about the changes in this PR.
--
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] manupa-arm commented on pull request #11453: [microNPU] Add transform matrices and part matcher to identity op
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on PR #11453:
URL: https://github.com/apache/tvm/pull/11453#issuecomment-1140885317
Hi @NicolaLancellotti ,
It would be great to rebase this PR on top of https://github.com/apache/tvm/pull/11410/files and see if the unit tests that I have disabled in that PR works -- so we can have higher confidence about the changes in this PR.
--
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 diff in pull request #11453: [microNPU] Add transform matrices and part matcher to identity op
Posted by GitBox <gi...@apache.org>.
NicolaLancellotti commented on code in PR #11453:
URL: https://github.com/apache/tvm/pull/11453#discussion_r882847066
##########
python/tvm/relay/backend/contrib/ethosu/te/identity.py:
##########
@@ -76,7 +78,85 @@ def identity_compute(
name="ethosu_identity",
attrs=id_attrs,
)
+ length = len(ifm.shape)
+ ifm_matrix = np.identity(length + 1)
+ offset = np.zeros(length, dtype="int64")
+ ifm_propagator = Propagator(
+ ifm_matrix,
+ offset.tolist(),
+ )
+ propagator_attrs = {
+ "ifm_propagator": ifm_propagator,
+ }
+ return write_compute(identity, ofm_zero_point, ofm_scale, attrs=propagator_attrs)
+
+
+@register_matcher
+def match_ethosu_identity(output_tensor, device_config):
+ """Match a Tensor Expression corresponding to an NPU identity.
- dmaed_ofm = write_compute(identity, ofm_zero_point, ofm_scale)
+ If the Tensor Expression matches, an EthosuPart will be created that models the
+ matched Tensor Expression. Otherwise, None will be returned.
- return dmaed_ofm
+ Parameters
+ ----------
+ output_tensor : tvm.te.Tensor
+ The tensor to attempt to match with.
+ device_config : EthosuDeviceConfig
+ Target device configuration
+
+ Returns
+ -------
+ Union[None, EthosuPart]
+ The created EthosuPart if there was a match, otherwise None.
+ """
+ write = output_tensor
+ if write.op.name != "ethosu_write":
+ return None
+ identity = write.op.input_tensors[0]
+ if identity.op.name != "ethosu_identity":
+ return None
+ read = identity.op.input_tensors[0]
+ if read.op.name != "ethosu_read":
+ return None
+
+ input_tensors = [
+ read.op.input_tensors[0],
+ ]
+ subgraph = TESubgraph(input_tensors, output_tensor)
+ propagators = [
+ write.op.attrs["ifm_propagator"],
+ ]
+ ifm_dtype = input_tensors[0].dtype
+ ofm_dtype = output_tensor.dtype
+
+ input_tensors_shape = input_tensors[0].shape
+ ifm_channels = int(input_tensors_shape[3] if len(input_tensors_shape) > 3 else 1)
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] lhutton1 commented on pull request #11453: [microNPU] Add transform matrices and part matcher to identity op
Posted by GitBox <gi...@apache.org>.
lhutton1 commented on PR #11453:
URL: https://github.com/apache/tvm/pull/11453#issuecomment-1143712503
Thanks @NicolaLancellotti, @ekalda, @manupa-arm!
--
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 diff in pull request #11453: [microNPU] Add transform matrices and part matcher to identity op
Posted by GitBox <gi...@apache.org>.
lhutton1 commented on code in PR #11453:
URL: https://github.com/apache/tvm/pull/11453#discussion_r885648438
##########
python/tvm/relay/backend/contrib/ethosu/te/identity.py:
##########
@@ -76,7 +78,85 @@ def identity_compute(
name="ethosu_identity",
attrs=id_attrs,
)
+ length = len(ifm.shape)
+ ifm_matrix = np.identity(length + 1)
+ offset = np.zeros(length, dtype="int64")
+ ifm_propagator = Propagator(
+ ifm_matrix,
+ offset.tolist(),
Review Comment:
Ah missed that, thanks!
--
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 diff in pull request #11453: [microNPU] Add transform matrices and part matcher to identity op
Posted by GitBox <gi...@apache.org>.
lhutton1 commented on code in PR #11453:
URL: https://github.com/apache/tvm/pull/11453#discussion_r886898788
##########
python/tvm/relay/backend/contrib/ethosu/te/identity.py:
##########
@@ -76,7 +78,85 @@ def identity_compute(
name="ethosu_identity",
attrs=id_attrs,
)
+ length = len(ifm.shape)
+ ifm_matrix = np.identity(length + 1)
+ offset = np.zeros(length, dtype="int64")
+ ifm_propagator = Propagator(
+ ifm_matrix,
+ offset.tolist(),
+ )
+ propagator_attrs = {
+ "ifm_propagator": ifm_propagator,
+ }
+ return write_compute(identity, ofm_zero_point, ofm_scale, attrs=propagator_attrs)
+
+
+@register_matcher
+def match_ethosu_identity(output_tensor, device_config):
+ """Match a Tensor Expression corresponding to an NPU identity.
- dmaed_ofm = write_compute(identity, ofm_zero_point, ofm_scale)
+ If the Tensor Expression matches, an EthosuPart will be created that models the
+ matched Tensor Expression. Otherwise, None will be returned.
- return dmaed_ofm
+ Parameters
+ ----------
+ output_tensor : tvm.te.Tensor
+ The tensor to attempt to match with.
+ device_config : EthosuDeviceConfig
+ Target device configuration
+
+ Returns
+ -------
+ Union[None, EthosuPart]
+ The created EthosuPart if there was a match, otherwise None.
+ """
+ write = output_tensor
+ if write.op.name != "ethosu_write":
+ return None
+ identity = write.op.input_tensors[0]
+ if identity.op.name != "ethosu_identity":
+ return None
+ read = identity.op.input_tensors[0]
+ if read.op.name != "ethosu_read":
+ return None
+
+ input_tensors = [
+ read.op.input_tensors[0],
+ ]
+ subgraph = TESubgraph(input_tensors, output_tensor)
+ propagators = [
+ write.op.attrs["ifm_propagator"],
+ ]
+ ifm_dtype = input_tensors[0].dtype
+ ofm_dtype = output_tensor.dtype
+
+ input_tensors_shape = input_tensors[0].shape
+ length = len(input_tensors_shape)
+ channels = int(input_tensors_shape[length - 1]) if length >= 3 else 1
+
+ subkernels = len(device_config.get_kernel_steps(identity.op.name, 1, 1, ifm_dtype))
+
+ input_layout = output_layout = "NHWC"
Review Comment:
Ah, apologies, I should have said that a message would have been helpful alongside the assert as well. Lets take it in a follow up :)
--
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 diff in pull request #11453: [microNPU] Add transform matrices and part matcher to identity op
Posted by GitBox <gi...@apache.org>.
lhutton1 commented on code in PR #11453:
URL: https://github.com/apache/tvm/pull/11453#discussion_r885606921
##########
python/tvm/relay/backend/contrib/ethosu/te/identity.py:
##########
@@ -76,7 +78,85 @@ def identity_compute(
name="ethosu_identity",
attrs=id_attrs,
)
+ length = len(ifm.shape)
+ ifm_matrix = np.identity(length + 1)
+ offset = np.zeros(length, dtype="int64")
+ ifm_propagator = Propagator(
+ ifm_matrix,
+ offset.tolist(),
+ )
+ propagator_attrs = {
+ "ifm_propagator": ifm_propagator,
+ }
+ return write_compute(identity, ofm_zero_point, ofm_scale, attrs=propagator_attrs)
+
+
+@register_matcher
+def match_ethosu_identity(output_tensor, device_config):
+ """Match a Tensor Expression corresponding to an NPU identity.
- dmaed_ofm = write_compute(identity, ofm_zero_point, ofm_scale)
+ If the Tensor Expression matches, an EthosuPart will be created that models the
+ matched Tensor Expression. Otherwise, None will be returned.
- return dmaed_ofm
+ Parameters
+ ----------
+ output_tensor : tvm.te.Tensor
+ The tensor to attempt to match with.
+ device_config : EthosuDeviceConfig
+ Target device configuration
+
+ Returns
+ -------
+ Union[None, EthosuPart]
+ The created EthosuPart if there was a match, otherwise None.
+ """
+ write = output_tensor
+ if write.op.name != "ethosu_write":
+ return None
+ identity = write.op.input_tensors[0]
+ if identity.op.name != "ethosu_identity":
+ return None
+ read = identity.op.input_tensors[0]
+ if read.op.name != "ethosu_read":
+ return None
+
+ input_tensors = [
+ read.op.input_tensors[0],
+ ]
+ subgraph = TESubgraph(input_tensors, output_tensor)
+ propagators = [
+ write.op.attrs["ifm_propagator"],
+ ]
+ ifm_dtype = input_tensors[0].dtype
+ ofm_dtype = output_tensor.dtype
+
+ input_tensors_shape = input_tensors[0].shape
+ length = len(input_tensors_shape)
+ channels = int(input_tensors_shape[length - 1]) if length >= 3 else 1
+
+ subkernels = len(device_config.get_kernel_steps(identity.op.name, 1, 1, ifm_dtype))
+
+ input_layout = output_layout = "NHWC"
Review Comment:
Perhaps we should assert that `len(input_tensors_shape) <= 4` if we don't support brick layout for identity
##########
python/tvm/relay/backend/contrib/ethosu/te/identity.py:
##########
@@ -76,7 +78,85 @@ def identity_compute(
name="ethosu_identity",
attrs=id_attrs,
)
+ length = len(ifm.shape)
+ ifm_matrix = np.identity(length + 1)
+ offset = np.zeros(length, dtype="int64")
+ ifm_propagator = Propagator(
+ ifm_matrix,
+ offset.tolist(),
Review Comment:
Seems a bit strange to need to use `.tolist()` for one of these and not the other, is it possible that its not needed?
--
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 diff in pull request #11453: [microNPU] Add transform matrices and part matcher to identity op
Posted by GitBox <gi...@apache.org>.
ekalda commented on code in PR #11453:
URL: https://github.com/apache/tvm/pull/11453#discussion_r882569177
##########
python/tvm/contrib/ethosu/cascader/device_config.py:
##########
@@ -48,9 +48,24 @@ def __init__(self, shape: List[int], layout="NHWC"):
self.width = int(shape[3])
self.depth = int(shape[2]) * int(shape[4])
else:
- self.height = int(shape[1])
- self.width = int(shape[2])
- self.depth = int(shape[3])
+ # identity layout is NHWC but the shape is not always 4
+ length = len(shape)
+ if length == 4:
+ self.height = int(shape[1])
+ self.width = int(shape[2])
+ self.depth = int(shape[3])
+ elif length == 3:
+ self.height = int(shape[1])
+ self.width = int(shape[2])
+ self.depth = 1
Review Comment:
I think?
```suggestion
elif length == 3:
self.height = int(shape[1])
self.width = int(shape[2])
self.depth = 1
```
```suggestion
elif length == 3:
self.height = int(shape[0])
self.width = int(shape[1])
self.depth = int(shape[2])
```
##########
python/tvm/contrib/ethosu/cascader/device_config.py:
##########
@@ -609,18 +624,19 @@ def _get_subkernel_propagator(
stride_w = int(op_attrs.get("stride_w", 1))
transform = ifm_propagator.transform
- if input_layout == "NHCWB16":
- transform[1][-1] = min(transform[1][-1], self._subkernel_limits[0] - stride_h)
- transform[3][-1] = min(transform[3][-1], self._subkernel_limits[1] - stride_w)
- else:
- transform[1][-1] = min(transform[1][-1], self._subkernel_limits[0] - stride_h)
- transform[2][-1] = min(transform[2][-1], self._subkernel_limits[1] - stride_w)
-
- if op_type in ("ethosu_pooling", "ethosu_depthwise_conv2d"):
- if output_layout == "NHCWB16" and input_layout == "NHWC":
- transform[3][-1] = depth
- elif output_layout == "NHCWB16" and input_layout == "NHCWB16":
- transform[2][-1] = 1 + ((depth - 1) // 16)
+ if op_type != "ethosu_identity":
Review Comment:
What happens if we don't hide that block behind the if statement?
##########
python/tvm/relay/backend/contrib/ethosu/te/identity.py:
##########
@@ -76,7 +78,85 @@ def identity_compute(
name="ethosu_identity",
attrs=id_attrs,
)
+ length = len(ifm.shape)
+ ifm_matrix = np.identity(length + 1)
+ offset = np.zeros(length, dtype="int64")
+ ifm_propagator = Propagator(
+ ifm_matrix,
+ offset.tolist(),
+ )
+ propagator_attrs = {
+ "ifm_propagator": ifm_propagator,
+ }
+ return write_compute(identity, ofm_zero_point, ofm_scale, attrs=propagator_attrs)
+
+
+@register_matcher
+def match_ethosu_identity(output_tensor, device_config):
+ """Match a Tensor Expression corresponding to an NPU identity.
- dmaed_ofm = write_compute(identity, ofm_zero_point, ofm_scale)
+ If the Tensor Expression matches, an EthosuPart will be created that models the
+ matched Tensor Expression. Otherwise, None will be returned.
- return dmaed_ofm
+ Parameters
+ ----------
+ output_tensor : tvm.te.Tensor
+ The tensor to attempt to match with.
+ device_config : EthosuDeviceConfig
+ Target device configuration
+
+ Returns
+ -------
+ Union[None, EthosuPart]
+ The created EthosuPart if there was a match, otherwise None.
+ """
+ write = output_tensor
+ if write.op.name != "ethosu_write":
+ return None
+ identity = write.op.input_tensors[0]
+ if identity.op.name != "ethosu_identity":
+ return None
+ read = identity.op.input_tensors[0]
+ if read.op.name != "ethosu_read":
+ return None
+
+ input_tensors = [
+ read.op.input_tensors[0],
+ ]
+ subgraph = TESubgraph(input_tensors, output_tensor)
+ propagators = [
+ write.op.attrs["ifm_propagator"],
+ ]
+ ifm_dtype = input_tensors[0].dtype
+ ofm_dtype = output_tensor.dtype
+
+ input_tensors_shape = input_tensors[0].shape
+ ifm_channels = int(input_tensors_shape[3] if len(input_tensors_shape) > 3 else 1)
Review Comment:
Since the 3D tensors most likely don't have the first dimension as 1, I think we should map them to (height, width, channels), e.g tensor with a shape (5, 6 ,7) should be mapped to (H=5, W=6, C=7)
##########
python/tvm/relay/backend/contrib/ethosu/te/identity.py:
##########
@@ -76,7 +78,85 @@ def identity_compute(
name="ethosu_identity",
attrs=id_attrs,
)
+ length = len(ifm.shape)
+ ifm_matrix = np.identity(length + 1)
+ offset = np.zeros(length, dtype="int64")
+ ifm_propagator = Propagator(
+ ifm_matrix,
+ offset.tolist(),
+ )
+ propagator_attrs = {
+ "ifm_propagator": ifm_propagator,
+ }
+ return write_compute(identity, ofm_zero_point, ofm_scale, attrs=propagator_attrs)
+
+
+@register_matcher
+def match_ethosu_identity(output_tensor, device_config):
+ """Match a Tensor Expression corresponding to an NPU identity.
- dmaed_ofm = write_compute(identity, ofm_zero_point, ofm_scale)
+ If the Tensor Expression matches, an EthosuPart will be created that models the
+ matched Tensor Expression. Otherwise, None will be returned.
- return dmaed_ofm
+ Parameters
+ ----------
+ output_tensor : tvm.te.Tensor
+ The tensor to attempt to match with.
+ device_config : EthosuDeviceConfig
+ Target device configuration
+
+ Returns
+ -------
+ Union[None, EthosuPart]
+ The created EthosuPart if there was a match, otherwise None.
+ """
+ write = output_tensor
+ if write.op.name != "ethosu_write":
+ return None
+ identity = write.op.input_tensors[0]
+ if identity.op.name != "ethosu_identity":
+ return None
+ read = identity.op.input_tensors[0]
+ if read.op.name != "ethosu_read":
+ return None
+
+ input_tensors = [
+ read.op.input_tensors[0],
+ ]
+ subgraph = TESubgraph(input_tensors, output_tensor)
+ propagators = [
+ write.op.attrs["ifm_propagator"],
+ ]
+ ifm_dtype = input_tensors[0].dtype
+ ofm_dtype = output_tensor.dtype
+
+ input_tensors_shape = input_tensors[0].shape
+ ifm_channels = int(input_tensors_shape[3] if len(input_tensors_shape) > 3 else 1)
+ ofm_channels = ifm_channels
Review Comment:
Nit: since `ifm_channels == ofm_channels`, maybe just use one variable (e.g. `channels`)
##########
python/tvm/contrib/ethosu/cascader/device_config.py:
##########
@@ -609,18 +624,19 @@ def _get_subkernel_propagator(
stride_w = int(op_attrs.get("stride_w", 1))
transform = ifm_propagator.transform
- if input_layout == "NHCWB16":
- transform[1][-1] = min(transform[1][-1], self._subkernel_limits[0] - stride_h)
- transform[3][-1] = min(transform[3][-1], self._subkernel_limits[1] - stride_w)
- else:
- transform[1][-1] = min(transform[1][-1], self._subkernel_limits[0] - stride_h)
- transform[2][-1] = min(transform[2][-1], self._subkernel_limits[1] - stride_w)
-
- if op_type in ("ethosu_pooling", "ethosu_depthwise_conv2d"):
- if output_layout == "NHCWB16" and input_layout == "NHWC":
- transform[3][-1] = depth
- elif output_layout == "NHCWB16" and input_layout == "NHCWB16":
- transform[2][-1] = 1 + ((depth - 1) // 16)
+ if op_type != "ethosu_identity":
Review Comment:
What happens if we don't hide that block behind the if statement?
##########
python/tvm/relay/backend/contrib/ethosu/te/identity.py:
##########
@@ -76,7 +78,85 @@ def identity_compute(
name="ethosu_identity",
attrs=id_attrs,
)
+ length = len(ifm.shape)
+ ifm_matrix = np.identity(length + 1)
+ offset = np.zeros(length, dtype="int64")
+ ifm_propagator = Propagator(
+ ifm_matrix,
+ offset.tolist(),
+ )
+ propagator_attrs = {
+ "ifm_propagator": ifm_propagator,
+ }
+ return write_compute(identity, ofm_zero_point, ofm_scale, attrs=propagator_attrs)
+
+
+@register_matcher
+def match_ethosu_identity(output_tensor, device_config):
+ """Match a Tensor Expression corresponding to an NPU identity.
- dmaed_ofm = write_compute(identity, ofm_zero_point, ofm_scale)
+ If the Tensor Expression matches, an EthosuPart will be created that models the
+ matched Tensor Expression. Otherwise, None will be returned.
- return dmaed_ofm
+ Parameters
+ ----------
+ output_tensor : tvm.te.Tensor
+ The tensor to attempt to match with.
+ device_config : EthosuDeviceConfig
+ Target device configuration
+
+ Returns
+ -------
+ Union[None, EthosuPart]
+ The created EthosuPart if there was a match, otherwise None.
+ """
+ write = output_tensor
+ if write.op.name != "ethosu_write":
+ return None
+ identity = write.op.input_tensors[0]
+ if identity.op.name != "ethosu_identity":
+ return None
+ read = identity.op.input_tensors[0]
+ if read.op.name != "ethosu_read":
+ return None
+
+ input_tensors = [
+ read.op.input_tensors[0],
+ ]
+ subgraph = TESubgraph(input_tensors, output_tensor)
+ propagators = [
+ write.op.attrs["ifm_propagator"],
+ ]
+ ifm_dtype = input_tensors[0].dtype
+ ofm_dtype = output_tensor.dtype
+
+ input_tensors_shape = input_tensors[0].shape
+ ifm_channels = int(input_tensors_shape[3] if len(input_tensors_shape) > 3 else 1)
Review Comment:
Since the 3D tensors most likely don't have the first dimension as 1, I think we should map them to (height, width, channels), e.g tensor with a shape (5, 6 ,7) should be mapped to (H=5, W=6, C=7)
##########
python/tvm/relay/backend/contrib/ethosu/te/identity.py:
##########
@@ -76,7 +78,85 @@ def identity_compute(
name="ethosu_identity",
attrs=id_attrs,
)
+ length = len(ifm.shape)
+ ifm_matrix = np.identity(length + 1)
+ offset = np.zeros(length, dtype="int64")
+ ifm_propagator = Propagator(
+ ifm_matrix,
+ offset.tolist(),
+ )
+ propagator_attrs = {
+ "ifm_propagator": ifm_propagator,
+ }
+ return write_compute(identity, ofm_zero_point, ofm_scale, attrs=propagator_attrs)
+
+
+@register_matcher
+def match_ethosu_identity(output_tensor, device_config):
+ """Match a Tensor Expression corresponding to an NPU identity.
- dmaed_ofm = write_compute(identity, ofm_zero_point, ofm_scale)
+ If the Tensor Expression matches, an EthosuPart will be created that models the
+ matched Tensor Expression. Otherwise, None will be returned.
- return dmaed_ofm
+ Parameters
+ ----------
+ output_tensor : tvm.te.Tensor
+ The tensor to attempt to match with.
+ device_config : EthosuDeviceConfig
+ Target device configuration
+
+ Returns
+ -------
+ Union[None, EthosuPart]
+ The created EthosuPart if there was a match, otherwise None.
+ """
+ write = output_tensor
+ if write.op.name != "ethosu_write":
+ return None
+ identity = write.op.input_tensors[0]
+ if identity.op.name != "ethosu_identity":
+ return None
+ read = identity.op.input_tensors[0]
+ if read.op.name != "ethosu_read":
+ return None
+
+ input_tensors = [
+ read.op.input_tensors[0],
+ ]
+ subgraph = TESubgraph(input_tensors, output_tensor)
+ propagators = [
+ write.op.attrs["ifm_propagator"],
+ ]
+ ifm_dtype = input_tensors[0].dtype
+ ofm_dtype = output_tensor.dtype
+
+ input_tensors_shape = input_tensors[0].shape
+ ifm_channels = int(input_tensors_shape[3] if len(input_tensors_shape) > 3 else 1)
+ ofm_channels = ifm_channels
Review Comment:
Nit: since `ifm_channels == ofm_channels`, maybe just use one variable (e.g. `channels`)
--
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 diff in pull request #11453: [microNPU] Add transform matrices and part matcher to identity op
Posted by GitBox <gi...@apache.org>.
NicolaLancellotti commented on code in PR #11453:
URL: https://github.com/apache/tvm/pull/11453#discussion_r882841312
##########
python/tvm/contrib/ethosu/cascader/device_config.py:
##########
@@ -48,9 +48,24 @@ def __init__(self, shape: List[int], layout="NHWC"):
self.width = int(shape[3])
self.depth = int(shape[2]) * int(shape[4])
else:
- self.height = int(shape[1])
- self.width = int(shape[2])
- self.depth = int(shape[3])
+ # identity layout is NHWC but the shape is not always 4
+ length = len(shape)
+ if length == 4:
+ self.height = int(shape[1])
+ self.width = int(shape[2])
+ self.depth = int(shape[3])
+ elif length == 3:
+ self.height = int(shape[1])
+ self.width = int(shape[2])
+ self.depth = 1
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