Mousius commented on a change in pull request #9471:
URL: https://github.com/apache/tvm/pull/9471#discussion_r780336416
##########
File path: python/tvm/relay/backend/contrib/ethosu/te/convolution.py
##########
@@ -164,5 +168,126 @@ def conv2d_compute(
attrs=conv2d_attrs,
)
+ nhwc_to_nhcwb16 = [
+ [1, 0, 0, 0, 0],
+ [0, 1, 0, 0, 0],
+ [0, 0, 0, 1 / 16, 0],
+ [0, 0, 1, 0, 0],
+ [0, 0, 0, 0, 16],
+ [0, 0, 0, 0, 1],
+ ]
+ nhcwb16_to_nhwc = [
+ [1, 0, 0, 0, 0, 0],
+ [0, 1, 0, 0, 0, 0],
+ [0, 0, 0, 1, 0, 0],
+ [0, 0, 16, 0, 0, 0],
+ [0, 0, 0, 0, 0, 1],
+ ]
+ ifm_matrix = [
+ [1, 0, 0, 0, 0],
+ [0, stride_h, 0, 0, (dilated_kernel_h - stride_h)],
+ [0, 0, stride_w, 0, (dilated_kernel_w - stride_w)],
+ [0, 0, 0, 0, ifm_channels],
+ [0, 0, 0, 0, 1],
+ ]
+ weights_matrix = [
+ [0, 0, 0, 1, 0],
+ [0, 0, 0, 0, kernel_h],
+ [0, 0, 0, 0, kernel_w],
+ [0, 0, 0, 0, ifm_channels],
+ [0, 0, 0, 0, 1],
+ ]
+ bias_matrix = [
+ [0, 0, 0, 1, 0],
+ [0, 0, 0, 0, 10],
+ [0, 0, 0, 0, 1],
+ ]
+ if ofm_layout == "NHCWB16":
+ ifm_matrix = np.matmul(ifm_matrix, nhcwb16_to_nhwc).tolist()
+ weights_matrix = np.matmul(weights_matrix, nhcwb16_to_nhwc).tolist()
+ bias_matrix = np.matmul(bias_matrix, nhcwb16_to_nhwc).tolist()
+ if ifm_layout == "NHCWB16":
+ ifm_matrix = np.matmul(nhwc_to_nhcwb16, ifm_matrix).tolist()
+ ifm_propagator = Propagator(
+ ifm_matrix,
+ [0, -padding[0], -padding[1], 0]
+ if ifm_layout == "NHWC"
+ else [0, -padding[0], 0, -padding[1], 0],
+ )
+ weights_propagator = Propagator(
+ weights_matrix,
+ [0, 0, 0, 0],
+ )
+ bias_propagator = Propagator(
+ bias_matrix,
+ [0, 0],
+ )
+ propagator_attrs = {
+ "ifm_propagator": ifm_propagator,
+ "weights_propagator": weights_propagator,
+ "bias_propagator": bias_propagator,
+ }
+
# Compute operation for the OFM DMA pipeline
- return dma_ofm_compute(conv, ofm_layout, ofm_zero_point, ofm_scale,
ofm_channels)
+ dma_ofm = dma_ofm_compute(
+ conv, ofm_layout, ofm_zero_point, ofm_scale, ofm_channels,
attrs=propagator_attrs
+ )
+ return dma_ofm
+
+
+@register_matcher
+def match_ethosu_conv2d(output_tensor):
+ """Match a Tensor Expression corresponding to an NPU Conv2D.
+
+ If the Tensor Expression matches, an EthosuPart will be created that
models the
+ matched Tensor Expression. Otherwise, None will be returned.
+
+ Parameters
+ ----------
+ output_tensor : tvm.te.Tensor
+ The tensor to attempt to match with.
+
+ 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
+ convert_to_nhcwb16 = write.op.input_tensors[0]
+ if convert_to_nhcwb16.op.name != "ethosu_convert_to_nhcwb16":
+ return None
+ conv2d = convert_to_nhcwb16.op.input_tensors[0]
+ if conv2d.op.name != "ethosu_conv2d":
+ return None
+ pad = conv2d.op.input_tensors[0]
+ if pad.op.name != "ethosu_pad":
+ return None
+ convert_to_nhwc = pad.op.input_tensors[0]
+ if convert_to_nhwc.op.name != "ethosu_convert_to_nhwc":
+ return None
+ read = convert_to_nhwc.op.input_tensors[0]
+ if read.op.name != "ethosu_read":
+ return None
Review comment:
I can't see test cases for all of these clauses, so we're missing most
of the logic of this function? Can you add some?
##########
File path:
tests/python/contrib/test_ethosu/cascader/test_ethosu_conv2d_matcher.py
##########
@@ -0,0 +1,163 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+import pytest
+
+pytest.importorskip("ethosu.vela")
+
+from tvm import te
+import tvm.contrib.ethosu.cascader as cs
+from tvm.relay.backend.contrib.ethosu.te.convolution import
match_ethosu_conv2d, conv2d_compute
+
+import numpy as np
+
+
+def _make_matrices(kernel, stride, dilation, padding, ifm_channels,
ifm_layout, ofm_layout):
+ kernel_h, kernel_w = kernel
+ stride_h, stride_w = stride
+ dilation_h, dilation_w = dilation
+ dilated_kernel_h = (kernel_h - 1) * dilation_h + 1
+ dilated_kernel_w = (kernel_w - 1) * dilation_w + 1
+ nhwc_to_nhcwb16 = [
+ [1, 0, 0, 0, 0],
+ [0, 1, 0, 0, 0],
+ [0, 0, 0, 1 / 16, 0],
+ [0, 0, 1, 0, 0],
+ [0, 0, 0, 0, 16],
+ [0, 0, 0, 0, 1],
+ ]
+ nhcwb16_to_nhwc = [
+ [1, 0, 0, 0, 0, 0],
+ [0, 1, 0, 0, 0, 0],
+ [0, 0, 0, 1, 0, 0],
+ [0, 0, 16, 0, 0, 0],
+ [0, 0, 0, 0, 0, 1],
+ ]
+ ifm_matrix = [
+ [1, 0, 0, 0, 0],
+ [0, stride_h, 0, 0, (dilated_kernel_h - stride_h)],
+ [0, 0, stride_w, 0, (dilated_kernel_w - stride_w)],
+ [0, 0, 0, 0, ifm_channels],
+ [0, 0, 0, 0, 1],
+ ]
+ weight_matrix = [
+ [0, 0, 0, 1, 0],
+ [0, 0, 0, 0, kernel_h],
+ [0, 0, 0, 0, kernel_w],
+ [0, 0, 0, 0, ifm_channels],
+ [0, 0, 0, 0, 1],
+ ]
+ scale_bias_matrix = [
+ [0, 0, 0, 1, 0],
+ [0, 0, 0, 0, 10],
+ [0, 0, 0, 0, 1],
+ ]
+ if ofm_layout == "NHCWB16":
+ ifm_matrix = np.matmul(ifm_matrix, nhcwb16_to_nhwc).tolist()
+ weight_matrix = np.matmul(weight_matrix, nhcwb16_to_nhwc).tolist()
+ scale_bias_matrix = np.matmul(scale_bias_matrix,
nhcwb16_to_nhwc).tolist()
+ if ifm_layout == "NHCWB16":
+ ifm_matrix = np.matmul(nhwc_to_nhcwb16, ifm_matrix).tolist()
+
+ ifm_offset = (
+ [0, -padding[0], -padding[1], 0]
+ if ifm_layout == "NHWC"
+ else [0, -padding[0], 0, -padding[1], 0]
+ )
+ weight_offset = [0, 0, 0, 0]
+ scale_bias_offset = [0, 0]
+ return (
+ ifm_matrix,
+ ifm_offset,
+ weight_matrix,
+ weight_offset,
+ scale_bias_matrix,
+ scale_bias_offset,
+ )
+
+
[email protected]("kernel", [(3, 3), (2, 1), (3, 5)])
[email protected]("stride", [(1, 1), (2, 1), (3, 2)])
[email protected]("dilation", [(1, 1), (2, 1), (3, 2)])
[email protected]("padding", [(0, 0, 0, 0), (3, 2, 3, 2), (2, 1, 0, 1)])
[email protected]("ifm_channels", [8, 57])
[email protected]("ifm_layout", ["NHWC", "NHCWB16"])
[email protected]("ofm_layout", ["NHWC", "NHCWB16"])
+def test_ethosu_conv2d_matcher(
+ kernel, stride, dilation, padding, ifm_channels, ifm_layout, ofm_layout
+):
+ if ifm_layout == "NHWC":
+ ifm_shape = (1, 12, 15, ifm_channels)
+ else:
+ ifm_shape = (1, 12, 1 + ((ifm_channels - 1) // 16), 15, 16)
+ ofm_channels = 8
+ kernel_h, kernel_w = kernel
+ ifm = te.placeholder(ifm_shape, dtype="int8")
+ weight = te.placeholder((ofm_channels, kernel_h, kernel_w, ifm_channels),
dtype="int8")
+ scale_bias = te.placeholder((ofm_channels, 10), dtype="uint8")
+ lut = te.placeholder((), dtype="uint8")
+ out = conv2d_compute(
+ ifm=ifm,
+ weight=weight,
+ scale_bias=scale_bias,
+ lut=lut,
+ ifm_scale=1,
+ ifm_zero_point=0,
+ ofm_scale=1,
+ ofm_zero_point=0,
+ weight_zero_point=0,
+ strides=stride,
+ padding=padding,
+ dilation=dilation,
+ activation="NONE",
+ clip_min=0,
+ clip_max=0,
+ upscale="NONE",
+ rounding_mode="TFL",
+ ifm_layout=ifm_layout,
+ ofm_layout=ofm_layout,
+ )
+ (
+ ifm_transform,
+ ifm_offset,
+ weight_transform,
+ weight_offset,
+ scale_bias_transform,
+ scale_bias_offset,
+ ) = _make_matrices(
+ kernel,
+ stride,
+ dilation,
+ padding,
+ ifm_channels,
+ ifm_layout,
+ ofm_layout,
+ )
+
+ part = match_ethosu_conv2d(out)
+
+ assert isinstance(part, cs.EthosuPart)
+ assert len(part.propagators) == 3
Review comment:
Can we test the other properties of the `part` returned?
##########
File path: tests/python/contrib/test_ethosu/cascader/test_graph.py
##########
@@ -16,14 +16,9 @@
# under the License.
import pytest
-from tvm.contrib.ethosu.cascader import (
- StripeConfig,
- Propagator,
- Tensor,
- InlinePart,
- TESubgraph,
- CascaderGraph,
-)
+pytest.importorskip("ethosu.vela")
+
+import tvm.contrib.ethosu.cascader as cs
Review comment:
As much as I appreciate using my initials here, it'd be really great to
use something I could immediately associate with the cascader; potentially
`cascader` rather than `cs` ?
##########
File path: python/tvm/contrib/ethosu/cascader/graph.py
##########
@@ -168,3 +174,80 @@ def tensor_order(self):
@property
def part_order(self):
return list(self._part_order)
+
+
+def register_matcher(matcher):
+ """Register a match function to the frontend.
+
+ A match function takes a te.Tensor and checks whether it matches
+ a known operator/operator sequence. If it does, it returns a Part
+ which models the behaviour of that operator sequence. Otherwise,
+ it returns None.
+ """
+ REGISTERED_MATCHERS.append(matcher)
+ return matcher
+
+
+def create_cascader_graph(te_graph: TESubgraph, const_dict: Dict[int,
np.ndarray]) -> CascaderGraph:
+ """Create a CascaderGraph from a Tensor Expression graph and constant
dictionary.
+
+ Parameters
+ ----------
+ te_graph : TESubgraph
+ The Tensor Expression graph.
+ const_dict : Dict[int, np.ndarray]
+ The constant dictionary.
+
+ Returns
+ -------
+ CascaderGraph
+ The CascaderGraph.
+ """
+ tensor_map = {}
+
+ def _visit_tensor(tensor):
+ if tensor not in tensor_map:
+ is_const = False
+ # Logic to determine if the tensor is constant
+ if tensor in list(te_graph.inputs):
+ i = list(te_graph.inputs).index(tensor)
+ if i in const_dict:
+ is_const = True
+
+ # TODO(@mbaret): Calculate the compression ratio
+ plan_tensor = Tensor(
+ tensor.shape,
+ tensor.dtype,
+ is_constant=is_const,
+ )
+ tensor_map[tensor] = plan_tensor
+ if isinstance(tensor.op, te.PlaceholderOp) or tensor in
te_graph.inputs:
+ return
+
+ input_tensors = []
+ # Check whether any of the registered matchers match the current
tensor
+ for matcher in REGISTERED_MATCHERS:
+ part = matcher(tensor)
+ if part:
+ input_tensors = part.subgraph.input_tensors
+ break
+
+ assert part is not None, f"The tensor {tensor} doesn't match any
part."
Review comment:
Would be good to ensure this fires when we expect it to with
`pytest.raises`.
##########
File path: python/tvm/relay/backend/contrib/ethosu/te/inline.py
##########
@@ -0,0 +1,71 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+"""Tensor Expressions for operations that will be inlined"""
+import numpy as np # type: ignore
+
+from tvm.contrib.ethosu.cascader import TESubgraph, InlinePart, Propagator,
register_matcher
+
+
+INLINE_OPS = {"T_reshape", "T_strided_slice"}
+
+
+@register_matcher
+def match_ethosu_inline(output_tensor):
+ """Match a Tensor Expression corresponding to an operator that will be
inlined.
+
+ If the Tensor Expression matches, an InlinePart will be created that
models the
+ matched Tensor Expression. Otherwise, None will be returned. This matcher
is
+ naive and assumes nothing about the compute of the Tensor Expression.
Therefore,
+ the resulting InlinePart will have full-tensor dependencies (i.e. each
output
+ element depends on every input element).
+
+ Parameters
+ ----------
+ output_tensor : tvm.te.Tensor
+ The tensor to attempt to match with.
+
+ Returns
+ -------
+ Union[None, InlinePart]
+ The created InlinePart if there was a match, otherwise None.
+
+ """
+ if output_tensor.op.name not in INLINE_OPS:
+ return None
Review comment:
Quick test for this please!
##########
File path:
tests/python/contrib/test_ethosu/cascader/test_ethosu_inline_matcher.py
##########
@@ -0,0 +1,49 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+import pytest
+
+pytest.importorskip("ethosu.vela")
+
+from tvm import te
+from tvm.topi.transform import reshape
+import tvm.contrib.ethosu.cascader as cs
+from tvm.relay.backend.contrib.ethosu.te.inline import match_ethosu_inline
+
+
+def test_ethosu_inline_matcher():
+ ifm_shape = (2, 5, 6)
+ new_shape = (2, 30)
+ ifm = te.placeholder(ifm_shape, dtype="int8")
+ out = reshape(ifm, new_shape)
+ ifm_transform = [
+ [0, 0, ifm_shape[0]],
+ [0, 0, ifm_shape[1]],
+ [0, 0, ifm_shape[2]],
+ [0, 0, 1],
+ ]
+ ifm_offset = [0, 0, 0]
+
+ part = match_ethosu_inline(out)
+
+ assert isinstance(part, cs.InlinePart)
+ assert len(part.propagators) == 1
Review comment:
Can we test the other properties of the `part` ?
##########
File path:
tests/python/contrib/test_ethosu/cascader/test_ethosu_inline_matcher.py
##########
@@ -0,0 +1,49 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+import pytest
+
+pytest.importorskip("ethosu.vela")
+
+from tvm import te
+from tvm.topi.transform import reshape
+import tvm.contrib.ethosu.cascader as cs
+from tvm.relay.backend.contrib.ethosu.te.inline import match_ethosu_inline
+
+
+def test_ethosu_inline_matcher():
+ ifm_shape = (2, 5, 6)
+ new_shape = (2, 30)
+ ifm = te.placeholder(ifm_shape, dtype="int8")
+ out = reshape(ifm, new_shape)
+ ifm_transform = [
+ [0, 0, ifm_shape[0]],
+ [0, 0, ifm_shape[1]],
+ [0, 0, ifm_shape[2]],
+ [0, 0, 1],
+ ]
+ ifm_offset = [0, 0, 0]
+
+ part = match_ethosu_inline(out)
+
+ assert isinstance(part, cs.InlinePart)
+ assert len(part.propagators) == 1
+ assert part.propagators[0].transform == ifm_transform
+ assert part.propagators[0].offset == ifm_offset
Review comment:
It would be good to test at least a few cases just to ensure the
function doesn't always return the same thing :smile_cat:
--
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]
