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: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
