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/23 20:53:48 UTC
[GitHub] [tvm] jverma-quic opened a new pull request, #11417: Implement avg_pool2d slice op
jverma-quic opened a new pull request, #11417:
URL: https://github.com/apache/tvm/pull/11417
Thanks for contributing to TVM! Please refer to guideline https://tvm.apache.org/docs/contribute/ for useful information and tips. After the pull request is submitted, please request code reviews from [Reviewers](https://github.com/apache/incubator-tvm/blob/master/CONTRIBUTORS.md#reviewers) by @ them in the pull request thread.
--
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] mehrdadh commented on a diff in pull request #11417: [Hexagon] Implement avg_pool2d slice op
Posted by GitBox <gi...@apache.org>.
mehrdadh commented on code in PR #11417:
URL: https://github.com/apache/tvm/pull/11417#discussion_r892599108
##########
tests/python/contrib/test_hexagon/test_avg_pool2d_slice.py:
##########
@@ -0,0 +1,369 @@
+# Licensed to the Apache Software Foundation (ASF) under one
Review Comment:
please move this file to `tests/python/contrib/test_hexagon/topi`
--
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] jverma-quic commented on a diff in pull request #11417: [Hexagon] Implement avg_pool2d slice op
Posted by GitBox <gi...@apache.org>.
jverma-quic commented on code in PR #11417:
URL: https://github.com/apache/tvm/pull/11417#discussion_r883957361
##########
python/tvm/topi/hexagon/utils.py:
##########
@@ -0,0 +1,75 @@
+# 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.
+
+from tvm import te
+
+
+def n11c_1024c_2d(n, h, w, c):
+ return [n, h, w, c // 1024, te.AXIS_SEPARATOR, c % 1024]
+
+
+def n11c_1024c_1d(n, h, w, c):
+ return [n, h, w, c // 1024, c % 1024]
+
+
+def nhwc_8h2w32c2w_2d(n, h, w, c):
+ return [n, h // 8, w // 4, c // 32, te.AXIS_SEPARATOR, h % 8, (w % 4) // 2, c % 32, w % 2]
+
+
+def nhwc_8h2w32c2w_1d(n, h, w, c):
+ return [n, h // 8, w // 4, c // 32, h % 8, (w % 4) // 2, c % 32, w % 2]
Review Comment:
Maybe. Although, I would prefer the current names.
--
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] kparzysz-quic merged pull request #11417: [Hexagon] Implement avg_pool2d slice op
Posted by GitBox <gi...@apache.org>.
kparzysz-quic merged PR #11417:
URL: https://github.com/apache/tvm/pull/11417
--
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] jverma-quic commented on pull request #11417: [Hexagon] Implement avg_pool2d slice op
Posted by GitBox <gi...@apache.org>.
jverma-quic commented on PR #11417:
URL: https://github.com/apache/tvm/pull/11417#issuecomment-1149047949
Hi @Lunderberg and @cconvey, please let me know if there are any additional comments.
--
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] cconvey commented on a diff in pull request #11417: Implement avg_pool2d slice op
Posted by GitBox <gi...@apache.org>.
cconvey commented on code in PR #11417:
URL: https://github.com/apache/tvm/pull/11417#discussion_r882091482
##########
tests/python/contrib/test_hexagon/infrastructure.py:
##########
@@ -228,3 +227,18 @@ def compute(n, ho, wo, ko, hi, wi, ki):
)
return output_shape, compute
+
+
+# Transpose and reshape numpy array according to the specified layout
+def transform_numpy(arr_np, layout):
Review Comment:
The function seems to assume that the supplied layout is NHWC. Is that a safe assumption for all expected uses of the function?
If no, then should we put `nhwc` into the function name, or perhaps change its argument list to something like `(arr_np, current_layout, new_layout)`?
--
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] jverma-quic commented on a diff in pull request #11417: [Hexagon] Implement avg_pool2d slice op
Posted by GitBox <gi...@apache.org>.
jverma-quic commented on code in PR #11417:
URL: https://github.com/apache/tvm/pull/11417#discussion_r883864069
##########
tests/python/contrib/test_hexagon/infrastructure.py:
##########
@@ -228,3 +227,18 @@ def compute(n, ho, wo, ko, hi, wi, ki):
)
return output_shape, compute
+
+
+# Transpose and reshape numpy array according to the specified layout
+def transform_numpy(arr_np, layout):
Review Comment:
> The function seems to assume that the supplied layout is NHWC. Is that a safe assumption for all expected uses of the function?
>
> If no, then should we put `nhwc` into the function name, or perhaps change its argument list to something like `(arr_np, current_layout, new_layout)`?
You're right that this function is making an assumption about the supplied layout which can transform the input incorrectly. I will include the current_layout as an argument. Thanks for the suggestion.
--
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] mehrdadh commented on a diff in pull request #11417: [Hexagon] Implement avg_pool2d slice op
Posted by GitBox <gi...@apache.org>.
mehrdadh commented on code in PR #11417:
URL: https://github.com/apache/tvm/pull/11417#discussion_r898195561
##########
tests/python/contrib/test_hexagon/topi/test_avg_pool2d_slice.py:
##########
@@ -0,0 +1,369 @@
+# 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
+import numpy as np
+
+from tvm import te, topi
+
+import tvm.testing
+from tvm.topi import testing
+from tvm.contrib.hexagon.build import HexagonLauncher
+import tvm.topi.hexagon.slice_ops as sl
+from ..infrastructure import allocate_hexagon_array, transform_numpy
+
+
+input_layout = tvm.testing.parameter(
+ "nhwc-8h2w32c2w-2d",
+)
+
+
+@tvm.testing.fixture
+def input_np(input_shape, dtype):
+ return np.random.random(input_shape).astype(dtype)
+
+
+@tvm.testing.fixture
+def transformed_expected_output_np(expected_output_np, output_layout):
+ return transform_numpy(expected_output_np, "nhwc", output_layout)
+
+
+@tvm.testing.fixture
+def transformed_input_np_padded(input_np_padded, input_layout):
+ return transform_numpy(input_np_padded, "nhwc", input_layout)
+
+
+class TestAvgPool2dSlice:
+ # NOTE: input_layout is always assumed to be "nhwc-8h2w32c2w-2d"
+ (
+ output_shape,
+ kernel,
+ stride,
+ dilation,
+ padding,
+ ceil_mode,
+ count_include_pad,
+ output_layout,
+ dtype,
+ ) = tvm.testing.parameters(
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w-2d",
+ "float16",
+ ),
+ (
+ [1, 16, 16, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w-2d",
+ "float16",
+ ),
+ (
+ [1, 8, 8, 32],
+ [8, 8],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w-2d",
+ "float16",
+ ),
+ # Test non-one stride and dilation
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [2, 3],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w-2d",
+ "float16",
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [2, 2],
+ [2, 2],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w-2d",
+ "float16",
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [2, 2],
+ [2, 3],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w-2d",
+ "float16",
+ ),
+ # Test non-zero padding
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [1, 1, 1, 1],
+ False,
+ True,
+ "nhwc-8h2w32c2w-2d",
+ "float16",
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [1, 2, 3, 4],
+ False,
+ True,
+ "nhwc-8h2w32c2w-2d",
+ "float16",
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [1, 2, 3, 4],
+ False,
+ True,
+ "nhwc-8h2w32c2w-2d",
+ "float16",
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [3, 2],
+ [2, 3],
+ [1, 2, 3, 4],
+ False,
+ True,
+ "nhwc-8h2w32c2w-2d",
+ "float16",
+ ),
+ # Test n11c-1024c-2d layout which will require input and output to have different layout
+ (
+ [1, 1, 1, 2048],
+ [8, 8],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c-2d",
+ "float16",
+ ),
+ (
+ [1, 1, 1, 2048],
+ [6, 6],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c-2d",
+ "float16",
+ ),
+ (
+ [1, 1, 1, 2048],
+ [3, 3],
+ [2, 2],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c-2d",
+ "float16",
+ ),
+ (
+ [1, 1, 1, 2048],
+ [4, 4],
+ [2, 2],
+ [2, 3],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c-2d",
+ "float16",
+ ),
+ )
+
+ @tvm.testing.fixture
+ def expected_output_np(
+ self,
+ input_np,
+ kernel,
+ stride,
+ dilation,
+ padding,
+ ceil_mode,
+ count_include_pad,
+ ):
+ pad_before = padding[:2]
+ pad_after = padding[2:]
+ ref_np = tvm.topi.testing.poolnd_python(
+ input_np,
+ kernel,
+ stride,
+ dilation,
+ pad_before,
+ pad_after,
+ "avg", # pool_type
+ count_include_pad,
+ False, # ceil_mode,
+ layout="NHWC",
+ )
+ return ref_np
+
+ @tvm.testing.fixture
+ def input_shape(self, output_shape, kernel, padding, stride, dilation, output_layout):
+ # Input shape without any padding; 'ceil' is being ignored from calculation:
+ o_b, o_h, o_w, o_c = output_shape
+ d_h, d_w = dilation
+ s_h, s_w = stride
+ k_h, k_w = kernel
+ pad_before_h, pad_before_w = padding[:2]
+ pad_after_h, pad_after_w = padding[2:]
+
+ if output_layout == "n11c-1024c-2d":
+ assert (
+ pad_before_w == 0 and pad_after_w == 0 and pad_before_h == 0 and pad_after_h == 0
+ ), "Padding must be zero for n11c-1024c-2d layout"
+ assert o_h == 1 and o_w == 1, "Output height and width must be 1"
+
+ in_h = (o_h - 1) * s_h + d_h * (k_h - 1) + 1 - pad_before_h - pad_after_h
+ in_w = (o_w - 1) * s_w + d_w * (k_w - 1) + 1 - pad_before_w - pad_after_w
+
+ return [o_b, in_h, in_w, o_c]
+
+ @tvm.testing.fixture
+ def input_shape_padded(self, input_shape, padding, output_layout):
+ # Input shape is adjusted to account for 'padding'. Also, due to the physical
+ # layout of the buffer, height and width are adjusted so that they are a
+ # multiple of 8 and 4 respectively.
+ # NOTE: Input layout is always assumed to be nhwc-8h2w32c2w-2d.
+ pad_before_h, pad_before_w = padding[:2]
+ pad_after_h, pad_after_w = padding[2:]
+ padded_input_height = ((input_shape[1] + pad_before_h + pad_after_h + 7) // 8) * 8
+ padded_input_width = ((input_shape[2] + pad_before_w + pad_after_w + 3) // 4) * 4
+ return [input_shape[0], padded_input_height, padded_input_width, input_shape[3]]
+
+ @tvm.testing.fixture
+ def input_np_padded(self, input_np, input_shape, input_shape_padded, padding):
+ pad_before_h, pad_before_w = padding[:2]
+ pad_after_h = input_shape_padded[1] - input_shape[1] - pad_before_h
+ pad_after_w = input_shape_padded[2] - input_shape[2] - pad_before_w
+ input_padded = np.pad(
+ input_np,
+ ((0, 0), (pad_before_h, pad_after_h), (pad_before_w, pad_after_w), (0, 0)),
+ "constant",
+ )
+ return input_padded
+
+ @tvm.testing.requires_hexagon
+ def test_avg_pool2d_slice(
+ self,
+ stride,
+ kernel,
+ dtype,
+ dilation,
+ padding,
+ count_include_pad,
+ input_layout,
+ output_layout,
+ output_shape,
+ input_shape,
+ input_shape_padded,
+ input_np,
+ input_np_padded,
+ transformed_input_np_padded,
+ transformed_expected_output_np,
+ expected_output_np,
+ hexagon_session,
+ ):
+
+ target_hexagon = tvm.target.hexagon("v69")
+ A = te.placeholder(input_shape_padded, name="A", dtype=dtype)
+
+ M = sl.avg_pool2d_compute(A, output_shape, kernel, stride, dilation)
+
+ # tir schedule
+ tir_schedule = sl.avg_pool2d_STIR_schedule(M, A, output_layout, input_layout)
+ sch = tir_schedule.mod
+
+ input_axis_separator = [4]
+ if output_layout == "nhwc-8h2w32c2w-2d":
+ output_axis_separator = [4]
+ elif output_layout == "n11c-1024c-2d":
+ output_axis_separator = [4]
+ else:
+ raise RuntimeError(f"Unexpected layout '{output_layout}'")
+
+ with tvm.transform.PassContext(opt_level=3):
+ func = tvm.build(
+ sch,
+ [A, M],
+ tvm.target.Target(target_hexagon, host=target_hexagon),
+ name="avg_pool2d",
+ )
+
+ input_arr = allocate_hexagon_array(
+ hexagon_session.device,
+ data=transformed_input_np_padded,
+ axis_separators=input_axis_separator,
+ mem_scope="global.vtcm",
+ )
+ output_arr = allocate_hexagon_array(
+ hexagon_session.device,
+ transformed_expected_output_np.shape,
+ dtype,
+ axis_separators=output_axis_separator,
+ mem_scope="global.vtcm",
+ )
+
+ mod = hexagon_session.load_module(func)
+ mod(input_arr, output_arr)
+ b, h, w, c = output_shape
+ if output_layout == "nhwc-8h2w32c2w-2d":
+ output_np = output_arr.numpy().reshape([b, h // 8, w // 4, c // 32, 8, 2, 32, 2])
+ elif output_layout == "n11c-1024c-2d":
+ output_np = output_arr.numpy().reshape([b, 1, 1, c // 1024, 1024])
+ else:
+ raise RuntimeError(f"Unexpected layout '{output_layout}'")
+
+ np.testing.assert_allclose(output_np, transformed_expected_output_np, rtol=1e-3, atol=1e-3)
+
+
+if __name__ == "__main__":
+ sys.exit(pytest.main(sys.argv))
Review Comment:
Don't need to update the PR only for this, but if you're adding more commit please change this line to:
`tvm.testing.main()`
--
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] csullivan commented on a diff in pull request #11417: [Hexagon] Implement avg_pool2d slice op
Posted by GitBox <gi...@apache.org>.
csullivan commented on code in PR #11417:
URL: https://github.com/apache/tvm/pull/11417#discussion_r898139232
##########
python/tvm/topi/hexagon/slice_ops/avg_pool2d.py:
##########
@@ -0,0 +1,141 @@
+# 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.
+# pylint: disable=invalid-name, unused-variable, unused-argument, too-many-locals
+
+""" Compute and schedule for avg_pool2d slice op
+
+Please note the following assumptions made by the implementation:
+
+1) The input must be padded in advance to account for 'padding'. In addition,
+ both input and output must be padded as per the physical buffer layout.
+2) The current implementation assumes 'count_include_pad' to be 'True'. It can be
+ modified to support 'False' case but the element count for the pooling window
+ must be pre-computed and provided as an input to reduce the run-time overhead.
+3) 'padding' is ignored. It must be handled outside of the sliced op.
+4) Please note that this implementation will not work if the output includes any
+ physical layout related padding as it can result into out-of-bound access
+ for the input.
+"""
+
+from tvm import te
+from tvm import tir
+from ..utils import get_layout_transform_fn
+
+
+def validate_out_shape(out_shape, in_shape, kernel, stride, dilation):
+ """Validate output shape"""
+ _, oh, ow, _ = out_shape
+ _, ih, iw, _ = in_shape
+ kh, kw = kernel
+ sh, sw = stride
+ dh, dw = dilation
+ if ih < (oh - 1) * sh + dh * (kh - 1) + 1:
+ raise RuntimeError("Output height is too large")
+ if iw < (ow - 1) * sw + dw * (kw - 1) + 1:
+ raise RuntimeError("Output width is too large")
+
+
+def avg_pool2d_compute(A, out_shape, kernel, stride, dilation):
+ """avg_pool2d compute"""
+ kh, kw = kernel
+ rh = te.reduce_axis((0, kh), name="rh")
+ rw = te.reduce_axis((0, kw), name="rw")
+ ob, oh, ow, oc = out_shape
+ if isinstance(ob, int):
+ validate_out_shape(out_shape, A.shape, kernel, stride, dilation)
+
+ sh, sw = stride
+ dh, dw = dilation
+ InvArea = float(1) / (kh * kw)
+
+ Sum = te.compute(
+ out_shape,
+ lambda b, h, w, c: te.sum(
+ A[b, h * sh + dh * rh, w * sw + dw * rw, c].astype("float32"), axis=[rh, rw]
+ ),
+ name="sum",
+ )
+ Avg = te.compute(
+ out_shape, lambda b, h, w, c: (Sum[b, h, w, c] * InvArea).astype(A.dtype), name="avg"
+ )
+ return Avg
+
+
+def STIR_schedule_nhwc_8h2w32c2w(outs, ins, output_layout: str, input_layout: str):
+ """Schedule for input and output layout nhwc-8h2w32c2w"""
+ func = te.create_prim_func([ins, outs])
+ s = tir.Schedule(func)
+ Sum = s.get_block("sum")
+ Avg = s.get_block("avg")
+
+ input_transform_fn = get_layout_transform_fn(input_layout)
+ output_transform_fn = get_layout_transform_fn(output_layout)
+ s.transform_layout(Sum, ("read", 0), input_transform_fn)
+ s.transform_layout(Avg, ("write", 0), output_transform_fn)
Review Comment:
Not necessary for this PR but just sharing: there's a new API sugar for transform_layout that allows you to address the block and buffer by name, e.g.
```
sch.transform_layout(block="tensor", buffer="placeholder", index_map=_int8_nhwc_8h8w32c_map)
```
--
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] Lunderberg commented on a diff in pull request #11417: [Hexagon] Implement avg_pool2d slice op
Posted by GitBox <gi...@apache.org>.
Lunderberg commented on code in PR #11417:
URL: https://github.com/apache/tvm/pull/11417#discussion_r887332670
##########
tests/python/contrib/test_hexagon/infrastructure.py:
##########
@@ -228,3 +227,18 @@ def compute(n, ho, wo, ko, hi, wi, ki):
)
return output_shape, compute
+
+
+# Transpose and reshape numpy array according to the specified layout
+def transform_numpy(arr_np, layout):
+ if layout == "nhwc":
+ return arr_np
+ elif layout == "nhwc-8h2w32c2w":
+ N, H, W, C = arr_np.shape
+ return arr_np.reshape([N, H // 8, 8, W // 4, 2, 2, C // 32, 32]).transpose(0, 1, 3, 6, 2, 4, 7, 5)
+ elif layout == "n11c-1024c":
+ N, H, W, C = arr_np.shape
+ assert (H == 1 and W == 1), "The size of H and W must be 1!"
+ return arr_np.reshape([N, C//1024, 1024]).transpose(0, 1, 2)
Review Comment:
No problem. I'm thinking that at some point, it might be useful to have a utility that converts a numpy array according to a `IndexMap`, to avoid needing to write the same transformation in numpy semantics.
--
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] TejashShah commented on pull request #11417: [Hexagon] Implement avg_pool2d slice op
Posted by GitBox <gi...@apache.org>.
TejashShah commented on PR #11417:
URL: https://github.com/apache/tvm/pull/11417#issuecomment-1150560308
Cc @Lunderberg @csullivan
--
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] cconvey commented on a diff in pull request #11417: Implement avg_pool2d slice op
Posted by GitBox <gi...@apache.org>.
cconvey commented on code in PR #11417:
URL: https://github.com/apache/tvm/pull/11417#discussion_r882086669
##########
python/tvm/topi/hexagon/utils.py:
##########
@@ -0,0 +1,75 @@
+# 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.
+
+from tvm import te
+
+
+def n11c_1024c_2d(n, h, w, c):
+ return [n, h, w, c // 1024, te.AXIS_SEPARATOR, c % 1024]
+
+
+def n11c_1024c_1d(n, h, w, c):
+ return [n, h, w, c // 1024, c % 1024]
+
+
+def nhwc_8h2w32c2w_2d(n, h, w, c):
+ return [n, h // 8, w // 4, c // 32, te.AXIS_SEPARATOR, h % 8, (w % 4) // 2, c % 32, w % 2]
+
+
+def nhwc_8h2w32c2w_1d(n, h, w, c):
+ return [n, h // 8, w // 4, c // 32, h % 8, (w % 4) // 2, c % 32, w % 2]
Review Comment:
Would these functions' purpose be clearer if their names started with something like `get_shape_...` or `transform_layout_...`?
--
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] TejashShah commented on pull request #11417: Implement avg_pool2d slice op
Posted by GitBox <gi...@apache.org>.
TejashShah commented on PR #11417:
URL: https://github.com/apache/tvm/pull/11417#issuecomment-1135146350
@csullivan @Lunderberg @cconvey, please review these changes.
--
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] cconvey commented on a diff in pull request #11417: [Hexagon] Implement avg_pool2d slice op
Posted by GitBox <gi...@apache.org>.
cconvey commented on code in PR #11417:
URL: https://github.com/apache/tvm/pull/11417#discussion_r882707828
##########
tests/python/contrib/test_hexagon/test_avg_pool2d_slice.py:
##########
@@ -0,0 +1,434 @@
+# 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
+import numpy as np
+
+np.set_printoptions(threshold=np.inf)
+from tvm import te, topi
+from tvm.tir.stmt_functor import post_order_visit
+
+import tvm.testing
+from tvm.topi import testing
+from tvm.contrib.hexagon.build import HexagonLauncher
+import tvm.topi.hexagon.slice_ops as sl
+from .infrastructure import allocate_hexagon_array, transform_numpy
+
+
+input_layout = tvm.testing.parameter(
+ "nhwc-8h2w32c2w",
+)
+
+
+@tvm.testing.fixture
+def input_np(input_shape, dtype):
+ return np.random.random(input_shape).astype(dtype)
+
+
+@tvm.testing.fixture
+def transformed_expected_output_np(expected_output_np, output_layout):
+ return transform_numpy(expected_output_np, output_layout)
+
+
+@tvm.testing.fixture
+def transformed_input_np_padded(input_np_padded, input_layout):
+ return transform_numpy(input_np_padded, input_layout)
+
+
+class TestAvgPool2dSlice:
+ # NOTE: input_layout is always assumed to be "nhwc-8h2w32c2w"
+ (
+ output_shape,
+ kernel,
+ stride,
+ dilation,
+ padding,
+ ceil_mode,
+ count_include_pad,
+ output_layout,
+ dtype,
+ use_te_sched,
+ ) = tvm.testing.parameters(
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 16, 16, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [8, 8],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [1, 1],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ # Test non-one stride and dilation
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [2, 3],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [2, 2],
+ [2, 2],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [2, 2],
+ [2, 3],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ # Test non-zero padding
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [1, 1, 1, 1],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [1, 2, 3, 4],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [1, 2, 3, 4],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [3, 2],
+ [2, 3],
+ [1, 2, 3, 4],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ # Test n11c-1024c layout which will require input and output to have different layout
+ (
+ [1, 1, 1, 2048],
+ [8, 8],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ False,
+ ),
+ (
+ [1, 1, 1, 2048],
+ [6, 6],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ False,
+ ),
+ (
+ [1, 1, 1, 2048],
+ [3, 3],
+ [2, 2],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ False,
+ ),
+ (
+ [1, 1, 1, 2048],
+ [4, 4],
+ [2, 2],
+ [2, 3],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ False,
+ ),
+ # Use 'te' schedule
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ True,
+ ),
+ (
+ [1, 1, 1, 2048],
+ [8, 8],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ True,
+ ),
+ )
+
+ @tvm.testing.fixture
+ def expected_output_np(
+ self,
+ input_np,
+ kernel,
+ stride,
+ dilation,
+ padding,
+ ceil_mode,
+ count_include_pad,
+ ):
+ pad_before = padding[:2]
+ pad_after = padding[2:]
+ ref_np = tvm.topi.testing.poolnd_python(
+ input_np,
+ kernel,
+ stride,
+ dilation,
+ pad_before,
+ pad_after,
+ "avg", # pool_type
+ count_include_pad,
+ False, # ceil_mode,
+ layout="NHWC",
+ )
+ return ref_np
+
+ @tvm.testing.fixture
+ def input_shape(self, output_shape, kernel, padding, stride, dilation, output_layout):
+ # Input shape without crouton padding; 'ceil' is being ignored from calculation:
+ o_b, o_h, o_w, o_c = output_shape
+ d_h, d_w = dilation
+ s_h, s_w = stride
+ k_h, k_w = kernel
+ pad_before_h, pad_before_w = padding[:2]
+ pad_after_h, pad_after_w = padding[2:]
+
+ if output_layout == "n11c-1024c":
+ assert (
+ pad_before_w == 0 and pad_after_w == 0 and pad_before_h == 0 and pad_after_h == 0
+ ), "Padding must be zero for n11c-1024c layout!!"
+ assert o_h == 1 and o_w == 1, "Output height and width must be 1!"
+
+ in_h = (o_h - 1) * s_h + d_h * (k_h - 1) + 1 - pad_before_h - pad_after_h
+ in_w = (o_w - 1) * s_w + d_w * (k_w - 1) + 1 - pad_before_w - pad_after_w
+
+ return [o_b, in_h, in_w, o_c]
+
+ @tvm.testing.fixture
+ def input_shape_padded(self, input_shape, padding, output_layout):
+ # Input shape with regular and crouton padding.
+ # Input width and height are padded to a multiple of croutons.
+ # NOTE: Input layout is always assumed to be nhwc-8h2w32c2w. Only the output layout can be
+ # different.
+ pad_before_h, pad_before_w = padding[:2]
+ pad_after_h, pad_after_w = padding[2:]
+ padded_input_height = ((input_shape[1] + pad_before_h + pad_after_h + 7) // 8) * 8
+ padded_input_width = ((input_shape[2] + pad_before_w + pad_after_w + 3) // 4) * 4
+ return [input_shape[0], padded_input_height, padded_input_width, input_shape[3]]
+
+ @tvm.testing.fixture
+ def input_np_padded(self, input_np, input_shape, input_shape_padded, padding):
+ pad_before_h, pad_before_w = padding[:2]
+ pad_after_h = (
+ input_shape_padded[1] - input_shape[1] - pad_before_h
+ ) # pad_after for height with crouton padding
+ pad_after_w = (
+ input_shape_padded[2] - input_shape[2] - pad_before_w
+ ) # pad_after for width with crouton padding
+ input_padded = np.pad(
+ input_np,
+ ((0, 0), (pad_before_h, pad_after_h), (pad_before_w, pad_after_w), (0, 0)),
+ "constant",
+ )
+ return input_padded
+
+ @tvm.testing.requires_hexagon
+ def test_avg_pool2d_slice(
+ self,
+ use_te_sched,
+ stride,
+ kernel,
+ dtype,
+ dilation,
+ padding,
+ count_include_pad,
+ input_layout,
+ output_layout,
+ output_shape,
+ input_shape,
+ input_shape_padded,
+ input_np,
+ input_np_padded,
+ transformed_input_np_padded,
+ transformed_expected_output_np,
+ expected_output_np,
+ hexagon_session,
+ ):
+
+ target_hexagon = tvm.target.hexagon("v69")
Review Comment:
AFAIK, TVM's Hexagon CI has only been run on v68 code in the past. Is there any possibility that specifying v69 here will break CI?
--
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] Lunderberg commented on a diff in pull request #11417: Implement avg_pool2d slice op
Posted by GitBox <gi...@apache.org>.
Lunderberg commented on code in PR #11417:
URL: https://github.com/apache/tvm/pull/11417#discussion_r882050980
##########
python/tvm/topi/hexagon/utils.py:
##########
@@ -0,0 +1,75 @@
+# 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.
+
+from tvm import te
+
+
+def n11c_1024c_2d(n, h, w, c):
+ return [n, h, w, c // 1024, te.AXIS_SEPARATOR, c % 1024]
+
+
+def n11c_1024c_1d(n, h, w, c):
+ return [n, h, w, c // 1024, c % 1024]
+
+
+def nhwc_8h2w32c2w_2d(n, h, w, c):
+ return [n, h // 8, w // 4, c // 32, te.AXIS_SEPARATOR, h % 8, (w % 4) // 2, c % 32, w % 2]
+
+
+def nhwc_8h2w32c2w_1d(n, h, w, c):
+ return [n, h // 8, w // 4, c // 32, h % 8, (w % 4) // 2, c % 32, w % 2]
+
+
+def get_layout_transform_fn(layout):
+ if layout == "nhwc-8h2w32c2w-2d":
+ return nhwc_8h2w32c2w_2d
+ if layout == "nhwc-8h2w32c2w-1d":
+ return nhwc_8h2w32c2w_1d
+ elif layout == "n11c-1024c-2d":
+ return n11c_1024c_2d
+ elif layout == "n11c-1024c-1d":
+ return n11c_1024c_1d
+ else:
+ raise RuntimeError(f"Unexpected layout '{layout}'")
+
+
+def apply_transform(s, block, block_index: int, buffer_type: str, layout: str):
+ """Apply transform layout on a buffer
+
+ Parameters
+ ----------
+ s: Schedule
+ block : BlockRV
+ The block that accesses the target buffer
+ buffer_index: int
+ The index of the buffer in block's read or write region
+ buffer_type : str
+ Type of the buffer index, "read" or "write"
+ layout : str
+ Layout of the buffer
+ """
+ transform_fn = get_layout_transform_fn(layout)
+ if layout == "nhwc-8h2w32c2w-1d":
+ axis_separators = [4]
+ elif layout == "n11c-1024c-1d":
+ axis_separators = [2]
+ else:
+ raise RuntimeError(f"Unexpected layout '{layout}'")
+
+ s.transform_layout(block, block_index, buffer_type, transform_fn)
Review Comment:
FYI, after #11269 lands, the calling `layout_transform` will also handle the call to `set_axis_separators`, so this function may become simpler or empty.
##########
python/tvm/topi/hexagon/slice_ops/avg_pool2d.py:
##########
@@ -0,0 +1,198 @@
+# 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.
+
+from tvm.ir.module import IRModule
+from tvm import te
+from tvm import tir
+from tvm.script import tir as T
+from ..utils import apply_transform, get_layout_transform_fn
+
+
+# The slice op implementation for avg_pool2d makes serveral assumptions:
+# 1) Both input and output are a multiple of croutons, and the input is already
+# padded for a given output shape as per any crouton and non-crouton related
+# padding.
+# 2) The current implementation assumes 'count_include_pad' to be 'True'. It can
+# modified to support 'False' but the element count for the pooling window must
+# be pre-computed and provided as an input to reduce the run-time overhead.
+# 3) 'padding' is also ignored. It must be handled outside of the sliced op.
+# 4) Please note that this implementation will not work if the output was padded
+# for the croutons. Since we loop over the logical output shape, this can result
+# into out-of-bound access for the input.
+
+def avg_pool2d_compute(A, out_shape, kernel, stride, dilation):
+ kh, kw = kernel
+ rh = te.reduce_axis((0, kh), name="rh")
+ rw = te.reduce_axis((0, kw), name="rw")
+ ob, oh, ow, oc = out_shape
+ sh, sw = stride
+ dh, dw = dilation
+ Area = float(1) / (kh * kw)
+
+ Sum = te.compute(
+ out_shape,
+ lambda b, h, w, c: te.sum(
+ A[b, h * sh + dh * rh, w * sw + dw * rw, c].astype("float32"), axis=[rh, rw]
+ ),
+ name="sum",
+ )
+ Avg = te.compute(
+ out_shape, lambda b, h, w, c: (Sum[b, h, w, c] * Area).astype(A.dtype), name="avg"
Review Comment:
Nitpick: The name `Area` threw me a bit, as I initially thought `Area` should be the area of the kernel relative to a single value, rather than the area of a value relative to the kernel. Can we rename `Area` to either `InvArea` or `NumValues`?
##########
python/tvm/topi/hexagon/slice_ops/avg_pool2d.py:
##########
@@ -0,0 +1,198 @@
+# 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.
+
+from tvm.ir.module import IRModule
+from tvm import te
+from tvm import tir
+from tvm.script import tir as T
+from ..utils import apply_transform, get_layout_transform_fn
+
+
+# The slice op implementation for avg_pool2d makes serveral assumptions:
+# 1) Both input and output are a multiple of croutons, and the input is already
+# padded for a given output shape as per any crouton and non-crouton related
+# padding.
+# 2) The current implementation assumes 'count_include_pad' to be 'True'. It can
+# modified to support 'False' but the element count for the pooling window must
+# be pre-computed and provided as an input to reduce the run-time overhead.
+# 3) 'padding' is also ignored. It must be handled outside of the sliced op.
+# 4) Please note that this implementation will not work if the output was padded
+# for the croutons. Since we loop over the logical output shape, this can result
+# into out-of-bound access for the input.
+
+def avg_pool2d_compute(A, out_shape, kernel, stride, dilation):
+ kh, kw = kernel
+ rh = te.reduce_axis((0, kh), name="rh")
+ rw = te.reduce_axis((0, kw), name="rw")
+ ob, oh, ow, oc = out_shape
+ sh, sw = stride
+ dh, dw = dilation
+ Area = float(1) / (kh * kw)
+
Review Comment:
Can we add validation to the `out_shape`, at least for static shapes? Since the external handling of padding means that we can't compute `out_shape` from the other parameters, it would be good to validate that `out_shape` isn't too large.
##########
python/tvm/topi/hexagon/slice_ops/avg_pool2d.py:
##########
@@ -0,0 +1,198 @@
+# 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.
+
+from tvm.ir.module import IRModule
+from tvm import te
+from tvm import tir
+from tvm.script import tir as T
+from ..utils import apply_transform, get_layout_transform_fn
+
+
+# The slice op implementation for avg_pool2d makes serveral assumptions:
Review Comment:
Can this be moved from a comment to a docstring?
##########
tests/python/contrib/test_hexagon/infrastructure.py:
##########
@@ -228,3 +227,18 @@ def compute(n, ho, wo, ko, hi, wi, ki):
)
return output_shape, compute
+
+
+# Transpose and reshape numpy array according to the specified layout
+def transform_numpy(arr_np, layout):
+ if layout == "nhwc":
+ return arr_np
+ elif layout == "nhwc-8h2w32c2w":
+ N, H, W, C = arr_np.shape
+ return arr_np.reshape([N, H // 8, 8, W // 4, 2, 2, C // 32, 32]).transpose(0, 1, 3, 6, 2, 4, 7, 5)
+ elif layout == "n11c-1024c":
+ N, H, W, C = arr_np.shape
+ assert (H == 1 and W == 1), "The size of H and W must be 1!"
+ return arr_np.reshape([N, C//1024, 1024]).transpose(0, 1, 2)
Review Comment:
This doesn't agree with the definition given in `topi.hexagon.utils.n11c_1024c_1d`. Even though it results in the same flattened shape when `H` and `W` are 1, they should still be included in order to pass the shape validation inserted in `MakePackedAPI`. Also, the `.transpose` should only be required if the dimensions are being reordered. The reshape below should match the layout transform defined in `topi.hexagon.utils.n11c_1024c_1d`.
```python
return arr_np.reshape([N, 1, 1, C//1024, 1024])
```
##########
python/tvm/topi/hexagon/slice_ops/avg_pool2d.py:
##########
@@ -0,0 +1,198 @@
+# 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.
+
+from tvm.ir.module import IRModule
+from tvm import te
+from tvm import tir
+from tvm.script import tir as T
+from ..utils import apply_transform, get_layout_transform_fn
+
+
+# The slice op implementation for avg_pool2d makes serveral assumptions:
+# 1) Both input and output are a multiple of croutons, and the input is already
+# padded for a given output shape as per any crouton and non-crouton related
+# padding.
+# 2) The current implementation assumes 'count_include_pad' to be 'True'. It can
+# modified to support 'False' but the element count for the pooling window must
+# be pre-computed and provided as an input to reduce the run-time overhead.
+# 3) 'padding' is also ignored. It must be handled outside of the sliced op.
+# 4) Please note that this implementation will not work if the output was padded
+# for the croutons. Since we loop over the logical output shape, this can result
+# into out-of-bound access for the input.
+
+def avg_pool2d_compute(A, out_shape, kernel, stride, dilation):
+ kh, kw = kernel
+ rh = te.reduce_axis((0, kh), name="rh")
+ rw = te.reduce_axis((0, kw), name="rw")
+ ob, oh, ow, oc = out_shape
+ sh, sw = stride
+ dh, dw = dilation
+ Area = float(1) / (kh * kw)
+
+ Sum = te.compute(
+ out_shape,
+ lambda b, h, w, c: te.sum(
+ A[b, h * sh + dh * rh, w * sw + dw * rw, c].astype("float32"), axis=[rh, rw]
+ ),
+ name="sum",
+ )
+ Avg = te.compute(
+ out_shape, lambda b, h, w, c: (Sum[b, h, w, c] * Area).astype(A.dtype), name="avg"
+ )
+ return Avg
+
+
+# Schedule for input and output layout nhwc-8h2w32c2w
+def STIR_schedule_nhwc_8h2w32c2w(outs, ins, output_layout: str, input_layout: str):
+ func = te.create_prim_func([ins, outs])
+ s = tir.Schedule(func)
+ Sum = s.get_block("sum")
+ Avg = s.get_block("avg")
+
+ apply_transform(s, Sum, 0, "read", input_layout)
+ apply_transform(s, Avg, 0, "write", output_layout)
+
+ # Schedule 'Sum'
+ bn, bh, bw, bc, rx, ry = s.get_loops(Sum)
+ bho, bhi = s.split(bh, [None, 8])
+ bwo, bwi = s.split(bw, [None, 4])
+ bwio, bwii = s.split(bwi, [None, 2]) # Doesn't seem to be doing anything
+ bco, bci = s.split(bc, [None, 32])
+ s.reorder(bn, bho, bwo, bco, bhi, bwio, rx, ry, bci, bwii) # --- DOESN'T do anything
+ bci_wii = s.fuse(bci, bwii) # --- DOESN'T do anything
Review Comment:
Same question here, after fusing I see extents `T.grid(1, 1, 2, 1, 8, 2, 3, 3, 64)` and can't reproduce the lack of effect.
##########
python/tvm/topi/hexagon/slice_ops/avg_pool2d.py:
##########
@@ -0,0 +1,198 @@
+# 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.
+
+from tvm.ir.module import IRModule
+from tvm import te
+from tvm import tir
+from tvm.script import tir as T
+from ..utils import apply_transform, get_layout_transform_fn
+
+
+# The slice op implementation for avg_pool2d makes serveral assumptions:
+# 1) Both input and output are a multiple of croutons, and the input is already
+# padded for a given output shape as per any crouton and non-crouton related
+# padding.
+# 2) The current implementation assumes 'count_include_pad' to be 'True'. It can
+# modified to support 'False' but the element count for the pooling window must
+# be pre-computed and provided as an input to reduce the run-time overhead.
+# 3) 'padding' is also ignored. It must be handled outside of the sliced op.
+# 4) Please note that this implementation will not work if the output was padded
+# for the croutons. Since we loop over the logical output shape, this can result
+# into out-of-bound access for the input.
+
+def avg_pool2d_compute(A, out_shape, kernel, stride, dilation):
+ kh, kw = kernel
+ rh = te.reduce_axis((0, kh), name="rh")
+ rw = te.reduce_axis((0, kw), name="rw")
+ ob, oh, ow, oc = out_shape
+ sh, sw = stride
+ dh, dw = dilation
+ Area = float(1) / (kh * kw)
+
+ Sum = te.compute(
+ out_shape,
+ lambda b, h, w, c: te.sum(
+ A[b, h * sh + dh * rh, w * sw + dw * rw, c].astype("float32"), axis=[rh, rw]
+ ),
+ name="sum",
+ )
+ Avg = te.compute(
+ out_shape, lambda b, h, w, c: (Sum[b, h, w, c] * Area).astype(A.dtype), name="avg"
+ )
+ return Avg
+
+
+# Schedule for input and output layout nhwc-8h2w32c2w
+def STIR_schedule_nhwc_8h2w32c2w(outs, ins, output_layout: str, input_layout: str):
+ func = te.create_prim_func([ins, outs])
+ s = tir.Schedule(func)
+ Sum = s.get_block("sum")
+ Avg = s.get_block("avg")
+
+ apply_transform(s, Sum, 0, "read", input_layout)
+ apply_transform(s, Avg, 0, "write", output_layout)
+
+ # Schedule 'Sum'
+ bn, bh, bw, bc, rx, ry = s.get_loops(Sum)
+ bho, bhi = s.split(bh, [None, 8])
+ bwo, bwi = s.split(bw, [None, 4])
+ bwio, bwii = s.split(bwi, [None, 2]) # Doesn't seem to be doing anything
+ bco, bci = s.split(bc, [None, 32])
+ s.reorder(bn, bho, bwo, bco, bhi, bwio, rx, ry, bci, bwii) # --- DOESN'T do anything
Review Comment:
What do you have before and after these lines? Running the test case `test_avg_pool2d_slice.py::TestAvgPool2dSlice::test_avg_pool2d_slice[nhwc-8h2w32c2w-False-str
ide0-kernel0-float16-dilation0-padding0-True-nhwc-8h2w32c2w-output_shape0-False]` and using `print(s.mod.script())`, I can see the loopnest before this line to have extents `T.grid(1, 1, 8, 2, 2, 2, 1, 32, 3, 3)` and afterward to have extents `T.grid(1, 1, 2, 1, 8, 2, 3, 3, 32, 2)`, so it does look like the reorder is having an effect.
--
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] cconvey commented on a diff in pull request #11417: [Hexagon] Implement avg_pool2d slice op
Posted by GitBox <gi...@apache.org>.
cconvey commented on code in PR #11417:
URL: https://github.com/apache/tvm/pull/11417#discussion_r882707828
##########
tests/python/contrib/test_hexagon/test_avg_pool2d_slice.py:
##########
@@ -0,0 +1,434 @@
+# 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
+import numpy as np
+
+np.set_printoptions(threshold=np.inf)
+from tvm import te, topi
+from tvm.tir.stmt_functor import post_order_visit
+
+import tvm.testing
+from tvm.topi import testing
+from tvm.contrib.hexagon.build import HexagonLauncher
+import tvm.topi.hexagon.slice_ops as sl
+from .infrastructure import allocate_hexagon_array, transform_numpy
+
+
+input_layout = tvm.testing.parameter(
+ "nhwc-8h2w32c2w",
+)
+
+
+@tvm.testing.fixture
+def input_np(input_shape, dtype):
+ return np.random.random(input_shape).astype(dtype)
+
+
+@tvm.testing.fixture
+def transformed_expected_output_np(expected_output_np, output_layout):
+ return transform_numpy(expected_output_np, output_layout)
+
+
+@tvm.testing.fixture
+def transformed_input_np_padded(input_np_padded, input_layout):
+ return transform_numpy(input_np_padded, input_layout)
+
+
+class TestAvgPool2dSlice:
+ # NOTE: input_layout is always assumed to be "nhwc-8h2w32c2w"
+ (
+ output_shape,
+ kernel,
+ stride,
+ dilation,
+ padding,
+ ceil_mode,
+ count_include_pad,
+ output_layout,
+ dtype,
+ use_te_sched,
+ ) = tvm.testing.parameters(
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 16, 16, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [8, 8],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [1, 1],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ # Test non-one stride and dilation
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [2, 3],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [2, 2],
+ [2, 2],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [2, 2],
+ [2, 3],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ # Test non-zero padding
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [1, 1, 1, 1],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [1, 2, 3, 4],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [1, 2, 3, 4],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [3, 2],
+ [2, 3],
+ [1, 2, 3, 4],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ # Test n11c-1024c layout which will require input and output to have different layout
+ (
+ [1, 1, 1, 2048],
+ [8, 8],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ False,
+ ),
+ (
+ [1, 1, 1, 2048],
+ [6, 6],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ False,
+ ),
+ (
+ [1, 1, 1, 2048],
+ [3, 3],
+ [2, 2],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ False,
+ ),
+ (
+ [1, 1, 1, 2048],
+ [4, 4],
+ [2, 2],
+ [2, 3],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ False,
+ ),
+ # Use 'te' schedule
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ True,
+ ),
+ (
+ [1, 1, 1, 2048],
+ [8, 8],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ True,
+ ),
+ )
+
+ @tvm.testing.fixture
+ def expected_output_np(
+ self,
+ input_np,
+ kernel,
+ stride,
+ dilation,
+ padding,
+ ceil_mode,
+ count_include_pad,
+ ):
+ pad_before = padding[:2]
+ pad_after = padding[2:]
+ ref_np = tvm.topi.testing.poolnd_python(
+ input_np,
+ kernel,
+ stride,
+ dilation,
+ pad_before,
+ pad_after,
+ "avg", # pool_type
+ count_include_pad,
+ False, # ceil_mode,
+ layout="NHWC",
+ )
+ return ref_np
+
+ @tvm.testing.fixture
+ def input_shape(self, output_shape, kernel, padding, stride, dilation, output_layout):
+ # Input shape without crouton padding; 'ceil' is being ignored from calculation:
+ o_b, o_h, o_w, o_c = output_shape
+ d_h, d_w = dilation
+ s_h, s_w = stride
+ k_h, k_w = kernel
+ pad_before_h, pad_before_w = padding[:2]
+ pad_after_h, pad_after_w = padding[2:]
+
+ if output_layout == "n11c-1024c":
+ assert (
+ pad_before_w == 0 and pad_after_w == 0 and pad_before_h == 0 and pad_after_h == 0
+ ), "Padding must be zero for n11c-1024c layout!!"
+ assert o_h == 1 and o_w == 1, "Output height and width must be 1!"
+
+ in_h = (o_h - 1) * s_h + d_h * (k_h - 1) + 1 - pad_before_h - pad_after_h
+ in_w = (o_w - 1) * s_w + d_w * (k_w - 1) + 1 - pad_before_w - pad_after_w
+
+ return [o_b, in_h, in_w, o_c]
+
+ @tvm.testing.fixture
+ def input_shape_padded(self, input_shape, padding, output_layout):
+ # Input shape with regular and crouton padding.
+ # Input width and height are padded to a multiple of croutons.
+ # NOTE: Input layout is always assumed to be nhwc-8h2w32c2w. Only the output layout can be
+ # different.
+ pad_before_h, pad_before_w = padding[:2]
+ pad_after_h, pad_after_w = padding[2:]
+ padded_input_height = ((input_shape[1] + pad_before_h + pad_after_h + 7) // 8) * 8
+ padded_input_width = ((input_shape[2] + pad_before_w + pad_after_w + 3) // 4) * 4
+ return [input_shape[0], padded_input_height, padded_input_width, input_shape[3]]
+
+ @tvm.testing.fixture
+ def input_np_padded(self, input_np, input_shape, input_shape_padded, padding):
+ pad_before_h, pad_before_w = padding[:2]
+ pad_after_h = (
+ input_shape_padded[1] - input_shape[1] - pad_before_h
+ ) # pad_after for height with crouton padding
+ pad_after_w = (
+ input_shape_padded[2] - input_shape[2] - pad_before_w
+ ) # pad_after for width with crouton padding
+ input_padded = np.pad(
+ input_np,
+ ((0, 0), (pad_before_h, pad_after_h), (pad_before_w, pad_after_w), (0, 0)),
+ "constant",
+ )
+ return input_padded
+
+ @tvm.testing.requires_hexagon
+ def test_avg_pool2d_slice(
+ self,
+ use_te_sched,
+ stride,
+ kernel,
+ dtype,
+ dilation,
+ padding,
+ count_include_pad,
+ input_layout,
+ output_layout,
+ output_shape,
+ input_shape,
+ input_shape_padded,
+ input_np,
+ input_np_padded,
+ transformed_input_np_padded,
+ transformed_expected_output_np,
+ expected_output_np,
+ hexagon_session,
+ ):
+
+ target_hexagon = tvm.target.hexagon("v69")
Review Comment:
AFAIK, TVM's Hexagon CI has only tested v68 code in the past. Is there any possibility that specifying v69 here will break CI or require some newer version of the Hexagon SDK?
--
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] jverma-quic commented on a diff in pull request #11417: [Hexagon] Implement avg_pool2d slice op
Posted by GitBox <gi...@apache.org>.
jverma-quic commented on code in PR #11417:
URL: https://github.com/apache/tvm/pull/11417#discussion_r883864069
##########
tests/python/contrib/test_hexagon/infrastructure.py:
##########
@@ -228,3 +227,18 @@ def compute(n, ho, wo, ko, hi, wi, ki):
)
return output_shape, compute
+
+
+# Transpose and reshape numpy array according to the specified layout
+def transform_numpy(arr_np, layout):
Review Comment:
> The function seems to assume that the supplied layout is NHWC. Is that a safe assumption for all expected uses of the function?
>
> If no, then should we put `nhwc` into the function name, or perhaps change its argument list to something like `(arr_np, current_layout, new_layout)`?
You're right that this function is making an assumption about the supplied layout which can transform the input incorrectly. I will include the current_layout as an argument. Thanks for the suggestion.
##########
python/tvm/topi/hexagon/slice_ops/avg_pool2d.py:
##########
@@ -0,0 +1,198 @@
+# 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.
+
+from tvm.ir.module import IRModule
+from tvm import te
+from tvm import tir
+from tvm.script import tir as T
+from ..utils import apply_transform, get_layout_transform_fn
+
+
+# The slice op implementation for avg_pool2d makes serveral assumptions:
+# 1) Both input and output are a multiple of croutons, and the input is already
+# padded for a given output shape as per any crouton and non-crouton related
+# padding.
+# 2) The current implementation assumes 'count_include_pad' to be 'True'. It can
+# modified to support 'False' but the element count for the pooling window must
+# be pre-computed and provided as an input to reduce the run-time overhead.
+# 3) 'padding' is also ignored. It must be handled outside of the sliced op.
+# 4) Please note that this implementation will not work if the output was padded
+# for the croutons. Since we loop over the logical output shape, this can result
+# into out-of-bound access for the input.
+
+def avg_pool2d_compute(A, out_shape, kernel, stride, dilation):
+ kh, kw = kernel
+ rh = te.reduce_axis((0, kh), name="rh")
+ rw = te.reduce_axis((0, kw), name="rw")
+ ob, oh, ow, oc = out_shape
+ sh, sw = stride
+ dh, dw = dilation
+ Area = float(1) / (kh * kw)
+
+ Sum = te.compute(
+ out_shape,
+ lambda b, h, w, c: te.sum(
+ A[b, h * sh + dh * rh, w * sw + dw * rw, c].astype("float32"), axis=[rh, rw]
+ ),
+ name="sum",
+ )
+ Avg = te.compute(
+ out_shape, lambda b, h, w, c: (Sum[b, h, w, c] * Area).astype(A.dtype), name="avg"
+ )
+ return Avg
+
+
+# Schedule for input and output layout nhwc-8h2w32c2w
+def STIR_schedule_nhwc_8h2w32c2w(outs, ins, output_layout: str, input_layout: str):
+ func = te.create_prim_func([ins, outs])
+ s = tir.Schedule(func)
+ Sum = s.get_block("sum")
+ Avg = s.get_block("avg")
+
+ apply_transform(s, Sum, 0, "read", input_layout)
+ apply_transform(s, Avg, 0, "write", output_layout)
+
+ # Schedule 'Sum'
+ bn, bh, bw, bc, rx, ry = s.get_loops(Sum)
+ bho, bhi = s.split(bh, [None, 8])
+ bwo, bwi = s.split(bw, [None, 4])
+ bwio, bwii = s.split(bwi, [None, 2]) # Doesn't seem to be doing anything
+ bco, bci = s.split(bc, [None, 32])
+ s.reorder(bn, bho, bwo, bco, bhi, bwio, rx, ry, bci, bwii) # --- DOESN'T do anything
Review Comment:
You're right. I do see the loops getting reordered after this line. However, when I print it again after s.compute_at(Sum, hi), I don't see the reordered/fused loopnest anymore.
--
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] jverma-quic commented on a diff in pull request #11417: [Hexagon] Implement avg_pool2d slice op
Posted by GitBox <gi...@apache.org>.
jverma-quic commented on code in PR #11417:
URL: https://github.com/apache/tvm/pull/11417#discussion_r898144950
##########
python/tvm/topi/hexagon/slice_ops/avg_pool2d.py:
##########
@@ -0,0 +1,141 @@
+# 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.
+# pylint: disable=invalid-name, unused-variable, unused-argument, too-many-locals
+
+""" Compute and schedule for avg_pool2d slice op
+
+Please note the following assumptions made by the implementation:
+
+1) The input must be padded in advance to account for 'padding'. In addition,
+ both input and output must be padded as per the physical buffer layout.
+2) The current implementation assumes 'count_include_pad' to be 'True'. It can be
+ modified to support 'False' case but the element count for the pooling window
+ must be pre-computed and provided as an input to reduce the run-time overhead.
+3) 'padding' is ignored. It must be handled outside of the sliced op.
+4) Please note that this implementation will not work if the output includes any
+ physical layout related padding as it can result into out-of-bound access
+ for the input.
+"""
+
+from tvm import te
+from tvm import tir
+from ..utils import get_layout_transform_fn
+
+
+def validate_out_shape(out_shape, in_shape, kernel, stride, dilation):
+ """Validate output shape"""
+ _, oh, ow, _ = out_shape
+ _, ih, iw, _ = in_shape
+ kh, kw = kernel
+ sh, sw = stride
+ dh, dw = dilation
+ if ih < (oh - 1) * sh + dh * (kh - 1) + 1:
+ raise RuntimeError("Output height is too large")
+ if iw < (ow - 1) * sw + dw * (kw - 1) + 1:
+ raise RuntimeError("Output width is too large")
+
+
+def avg_pool2d_compute(A, out_shape, kernel, stride, dilation):
+ """avg_pool2d compute"""
+ kh, kw = kernel
+ rh = te.reduce_axis((0, kh), name="rh")
+ rw = te.reduce_axis((0, kw), name="rw")
+ ob, oh, ow, oc = out_shape
+ if isinstance(ob, int):
+ validate_out_shape(out_shape, A.shape, kernel, stride, dilation)
+
+ sh, sw = stride
+ dh, dw = dilation
+ InvArea = float(1) / (kh * kw)
+
+ Sum = te.compute(
+ out_shape,
+ lambda b, h, w, c: te.sum(
+ A[b, h * sh + dh * rh, w * sw + dw * rw, c].astype("float32"), axis=[rh, rw]
+ ),
+ name="sum",
+ )
+ Avg = te.compute(
+ out_shape, lambda b, h, w, c: (Sum[b, h, w, c] * InvArea).astype(A.dtype), name="avg"
+ )
+ return Avg
+
+
+def STIR_schedule_nhwc_8h2w32c2w(outs, ins, output_layout: str, input_layout: str):
+ """Schedule for input and output layout nhwc-8h2w32c2w"""
+ func = te.create_prim_func([ins, outs])
+ s = tir.Schedule(func)
+ Sum = s.get_block("sum")
+ Avg = s.get_block("avg")
+
+ input_transform_fn = get_layout_transform_fn(input_layout)
+ output_transform_fn = get_layout_transform_fn(output_layout)
+ s.transform_layout(Sum, ("read", 0), input_transform_fn)
+ s.transform_layout(Avg, ("write", 0), output_transform_fn)
Review Comment:
Thanks for the suggestion, @csullivan! I have tried it in the past but couldn't get the new API sugar to work due to the intermediate compute in avgpool.
--
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] cconvey commented on a diff in pull request #11417: [Hexagon] Implement avg_pool2d slice op
Posted by GitBox <gi...@apache.org>.
cconvey commented on code in PR #11417:
URL: https://github.com/apache/tvm/pull/11417#discussion_r882707828
##########
tests/python/contrib/test_hexagon/test_avg_pool2d_slice.py:
##########
@@ -0,0 +1,434 @@
+# 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
+import numpy as np
+
+np.set_printoptions(threshold=np.inf)
+from tvm import te, topi
+from tvm.tir.stmt_functor import post_order_visit
+
+import tvm.testing
+from tvm.topi import testing
+from tvm.contrib.hexagon.build import HexagonLauncher
+import tvm.topi.hexagon.slice_ops as sl
+from .infrastructure import allocate_hexagon_array, transform_numpy
+
+
+input_layout = tvm.testing.parameter(
+ "nhwc-8h2w32c2w",
+)
+
+
+@tvm.testing.fixture
+def input_np(input_shape, dtype):
+ return np.random.random(input_shape).astype(dtype)
+
+
+@tvm.testing.fixture
+def transformed_expected_output_np(expected_output_np, output_layout):
+ return transform_numpy(expected_output_np, output_layout)
+
+
+@tvm.testing.fixture
+def transformed_input_np_padded(input_np_padded, input_layout):
+ return transform_numpy(input_np_padded, input_layout)
+
+
+class TestAvgPool2dSlice:
+ # NOTE: input_layout is always assumed to be "nhwc-8h2w32c2w"
+ (
+ output_shape,
+ kernel,
+ stride,
+ dilation,
+ padding,
+ ceil_mode,
+ count_include_pad,
+ output_layout,
+ dtype,
+ use_te_sched,
+ ) = tvm.testing.parameters(
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 16, 16, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [8, 8],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [1, 1],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ # Test non-one stride and dilation
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [2, 3],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [2, 2],
+ [2, 2],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [2, 2],
+ [2, 3],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ # Test non-zero padding
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [1, 1, 1, 1],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [1, 2, 3, 4],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [1, 2, 3, 4],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [3, 2],
+ [2, 3],
+ [1, 2, 3, 4],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ # Test n11c-1024c layout which will require input and output to have different layout
+ (
+ [1, 1, 1, 2048],
+ [8, 8],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ False,
+ ),
+ (
+ [1, 1, 1, 2048],
+ [6, 6],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ False,
+ ),
+ (
+ [1, 1, 1, 2048],
+ [3, 3],
+ [2, 2],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ False,
+ ),
+ (
+ [1, 1, 1, 2048],
+ [4, 4],
+ [2, 2],
+ [2, 3],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ False,
+ ),
+ # Use 'te' schedule
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ True,
+ ),
+ (
+ [1, 1, 1, 2048],
+ [8, 8],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ True,
+ ),
+ )
+
+ @tvm.testing.fixture
+ def expected_output_np(
+ self,
+ input_np,
+ kernel,
+ stride,
+ dilation,
+ padding,
+ ceil_mode,
+ count_include_pad,
+ ):
+ pad_before = padding[:2]
+ pad_after = padding[2:]
+ ref_np = tvm.topi.testing.poolnd_python(
+ input_np,
+ kernel,
+ stride,
+ dilation,
+ pad_before,
+ pad_after,
+ "avg", # pool_type
+ count_include_pad,
+ False, # ceil_mode,
+ layout="NHWC",
+ )
+ return ref_np
+
+ @tvm.testing.fixture
+ def input_shape(self, output_shape, kernel, padding, stride, dilation, output_layout):
+ # Input shape without crouton padding; 'ceil' is being ignored from calculation:
+ o_b, o_h, o_w, o_c = output_shape
+ d_h, d_w = dilation
+ s_h, s_w = stride
+ k_h, k_w = kernel
+ pad_before_h, pad_before_w = padding[:2]
+ pad_after_h, pad_after_w = padding[2:]
+
+ if output_layout == "n11c-1024c":
+ assert (
+ pad_before_w == 0 and pad_after_w == 0 and pad_before_h == 0 and pad_after_h == 0
+ ), "Padding must be zero for n11c-1024c layout!!"
+ assert o_h == 1 and o_w == 1, "Output height and width must be 1!"
+
+ in_h = (o_h - 1) * s_h + d_h * (k_h - 1) + 1 - pad_before_h - pad_after_h
+ in_w = (o_w - 1) * s_w + d_w * (k_w - 1) + 1 - pad_before_w - pad_after_w
+
+ return [o_b, in_h, in_w, o_c]
+
+ @tvm.testing.fixture
+ def input_shape_padded(self, input_shape, padding, output_layout):
+ # Input shape with regular and crouton padding.
+ # Input width and height are padded to a multiple of croutons.
+ # NOTE: Input layout is always assumed to be nhwc-8h2w32c2w. Only the output layout can be
+ # different.
+ pad_before_h, pad_before_w = padding[:2]
+ pad_after_h, pad_after_w = padding[2:]
+ padded_input_height = ((input_shape[1] + pad_before_h + pad_after_h + 7) // 8) * 8
+ padded_input_width = ((input_shape[2] + pad_before_w + pad_after_w + 3) // 4) * 4
+ return [input_shape[0], padded_input_height, padded_input_width, input_shape[3]]
+
+ @tvm.testing.fixture
+ def input_np_padded(self, input_np, input_shape, input_shape_padded, padding):
+ pad_before_h, pad_before_w = padding[:2]
+ pad_after_h = (
+ input_shape_padded[1] - input_shape[1] - pad_before_h
+ ) # pad_after for height with crouton padding
+ pad_after_w = (
+ input_shape_padded[2] - input_shape[2] - pad_before_w
+ ) # pad_after for width with crouton padding
+ input_padded = np.pad(
+ input_np,
+ ((0, 0), (pad_before_h, pad_after_h), (pad_before_w, pad_after_w), (0, 0)),
+ "constant",
+ )
+ return input_padded
+
+ @tvm.testing.requires_hexagon
+ def test_avg_pool2d_slice(
+ self,
+ use_te_sched,
+ stride,
+ kernel,
+ dtype,
+ dilation,
+ padding,
+ count_include_pad,
+ input_layout,
+ output_layout,
+ output_shape,
+ input_shape,
+ input_shape_padded,
+ input_np,
+ input_np_padded,
+ transformed_input_np_padded,
+ transformed_expected_output_np,
+ expected_output_np,
+ hexagon_session,
+ ):
+
+ target_hexagon = tvm.target.hexagon("v69")
Review Comment:
AFAIK, TVM's Hexagon CI has only been run on v68 code in the past. Is there any possibility that specifying v69 here will break CI or require some newer version of the Hexagon SDK?
--
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] jverma-quic commented on a diff in pull request #11417: [Hexagon] Implement avg_pool2d slice op
Posted by GitBox <gi...@apache.org>.
jverma-quic commented on code in PR #11417:
URL: https://github.com/apache/tvm/pull/11417#discussion_r883953282
##########
tests/python/contrib/test_hexagon/infrastructure.py:
##########
@@ -228,3 +227,18 @@ def compute(n, ho, wo, ko, hi, wi, ki):
)
return output_shape, compute
+
+
+# Transpose and reshape numpy array according to the specified layout
+def transform_numpy(arr_np, layout):
+ if layout == "nhwc":
+ return arr_np
+ elif layout == "nhwc-8h2w32c2w":
+ N, H, W, C = arr_np.shape
+ return arr_np.reshape([N, H // 8, 8, W // 4, 2, 2, C // 32, 32]).transpose(0, 1, 3, 6, 2, 4, 7, 5)
+ elif layout == "n11c-1024c":
+ N, H, W, C = arr_np.shape
+ assert (H == 1 and W == 1), "The size of H and W must be 1!"
+ return arr_np.reshape([N, C//1024, 1024]).transpose(0, 1, 2)
Review Comment:
Thanks for pointing it out.
##########
tests/python/contrib/test_hexagon/test_avg_pool2d_slice.py:
##########
@@ -0,0 +1,434 @@
+# 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
+import numpy as np
+
+np.set_printoptions(threshold=np.inf)
+from tvm import te, topi
+from tvm.tir.stmt_functor import post_order_visit
+
+import tvm.testing
+from tvm.topi import testing
+from tvm.contrib.hexagon.build import HexagonLauncher
+import tvm.topi.hexagon.slice_ops as sl
+from .infrastructure import allocate_hexagon_array, transform_numpy
+
+
+input_layout = tvm.testing.parameter(
+ "nhwc-8h2w32c2w",
+)
+
+
+@tvm.testing.fixture
+def input_np(input_shape, dtype):
+ return np.random.random(input_shape).astype(dtype)
+
+
+@tvm.testing.fixture
+def transformed_expected_output_np(expected_output_np, output_layout):
+ return transform_numpy(expected_output_np, output_layout)
+
+
+@tvm.testing.fixture
+def transformed_input_np_padded(input_np_padded, input_layout):
+ return transform_numpy(input_np_padded, input_layout)
+
+
+class TestAvgPool2dSlice:
+ # NOTE: input_layout is always assumed to be "nhwc-8h2w32c2w"
+ (
+ output_shape,
+ kernel,
+ stride,
+ dilation,
+ padding,
+ ceil_mode,
+ count_include_pad,
+ output_layout,
+ dtype,
+ use_te_sched,
+ ) = tvm.testing.parameters(
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 16, 16, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [8, 8],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [1, 1],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ # Test non-one stride and dilation
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [2, 3],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [2, 2],
+ [2, 2],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [2, 2],
+ [2, 3],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ # Test non-zero padding
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [1, 1, 1, 1],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [1, 2, 3, 4],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [1, 2, 3, 4],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [3, 2],
+ [2, 3],
+ [1, 2, 3, 4],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ # Test n11c-1024c layout which will require input and output to have different layout
+ (
+ [1, 1, 1, 2048],
+ [8, 8],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ False,
+ ),
+ (
+ [1, 1, 1, 2048],
+ [6, 6],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ False,
+ ),
+ (
+ [1, 1, 1, 2048],
+ [3, 3],
+ [2, 2],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ False,
+ ),
+ (
+ [1, 1, 1, 2048],
+ [4, 4],
+ [2, 2],
+ [2, 3],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ False,
+ ),
+ # Use 'te' schedule
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ True,
+ ),
+ (
+ [1, 1, 1, 2048],
+ [8, 8],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ True,
+ ),
+ )
+
+ @tvm.testing.fixture
+ def expected_output_np(
+ self,
+ input_np,
+ kernel,
+ stride,
+ dilation,
+ padding,
+ ceil_mode,
+ count_include_pad,
+ ):
+ pad_before = padding[:2]
+ pad_after = padding[2:]
+ ref_np = tvm.topi.testing.poolnd_python(
+ input_np,
+ kernel,
+ stride,
+ dilation,
+ pad_before,
+ pad_after,
+ "avg", # pool_type
+ count_include_pad,
+ False, # ceil_mode,
+ layout="NHWC",
+ )
+ return ref_np
+
+ @tvm.testing.fixture
+ def input_shape(self, output_shape, kernel, padding, stride, dilation, output_layout):
+ # Input shape without crouton padding; 'ceil' is being ignored from calculation:
+ o_b, o_h, o_w, o_c = output_shape
+ d_h, d_w = dilation
+ s_h, s_w = stride
+ k_h, k_w = kernel
+ pad_before_h, pad_before_w = padding[:2]
+ pad_after_h, pad_after_w = padding[2:]
+
+ if output_layout == "n11c-1024c":
+ assert (
+ pad_before_w == 0 and pad_after_w == 0 and pad_before_h == 0 and pad_after_h == 0
+ ), "Padding must be zero for n11c-1024c layout!!"
Review Comment:
Agreed. :)
--
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] jverma-quic commented on a diff in pull request #11417: [Hexagon] Implement avg_pool2d slice op
Posted by GitBox <gi...@apache.org>.
jverma-quic commented on code in PR #11417:
URL: https://github.com/apache/tvm/pull/11417#discussion_r883958719
##########
tests/python/contrib/test_hexagon/infrastructure.py:
##########
@@ -228,3 +227,18 @@ def compute(n, ho, wo, ko, hi, wi, ki):
)
return output_shape, compute
+
+
+# Transpose and reshape numpy array according to the specified layout
+def transform_numpy(arr_np, layout):
Review Comment:
You're right that this function is making an assumption about the supplied layout which can transform the input incorrectly. Thanks for the suggestion, @cconvey!
--
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] cconvey commented on a diff in pull request #11417: Implement avg_pool2d slice op
Posted by GitBox <gi...@apache.org>.
cconvey commented on code in PR #11417:
URL: https://github.com/apache/tvm/pull/11417#discussion_r882093310
##########
tests/python/contrib/test_hexagon/test_avg_pool2d_slice.py:
##########
@@ -0,0 +1,434 @@
+# 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
+import numpy as np
+
+np.set_printoptions(threshold=np.inf)
+from tvm import te, topi
+from tvm.tir.stmt_functor import post_order_visit
+
+import tvm.testing
+from tvm.topi import testing
+from tvm.contrib.hexagon.build import HexagonLauncher
+import tvm.topi.hexagon.slice_ops as sl
+from .infrastructure import allocate_hexagon_array, transform_numpy
+
+
+input_layout = tvm.testing.parameter(
+ "nhwc-8h2w32c2w",
+)
+
+
+@tvm.testing.fixture
+def input_np(input_shape, dtype):
+ return np.random.random(input_shape).astype(dtype)
+
+
+@tvm.testing.fixture
+def transformed_expected_output_np(expected_output_np, output_layout):
+ return transform_numpy(expected_output_np, output_layout)
+
+
+@tvm.testing.fixture
+def transformed_input_np_padded(input_np_padded, input_layout):
+ return transform_numpy(input_np_padded, input_layout)
+
+
+class TestAvgPool2dSlice:
+ # NOTE: input_layout is always assumed to be "nhwc-8h2w32c2w"
+ (
+ output_shape,
+ kernel,
+ stride,
+ dilation,
+ padding,
+ ceil_mode,
+ count_include_pad,
+ output_layout,
+ dtype,
+ use_te_sched,
+ ) = tvm.testing.parameters(
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 16, 16, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [8, 8],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [1, 1],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ # Test non-one stride and dilation
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [2, 3],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [2, 2],
+ [2, 2],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [2, 2],
+ [2, 3],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ # Test non-zero padding
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [1, 1, 1, 1],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [1, 2, 3, 4],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [1, 2, 3, 4],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [3, 2],
+ [2, 3],
+ [1, 2, 3, 4],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ # Test n11c-1024c layout which will require input and output to have different layout
+ (
+ [1, 1, 1, 2048],
+ [8, 8],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ False,
+ ),
+ (
+ [1, 1, 1, 2048],
+ [6, 6],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ False,
+ ),
+ (
+ [1, 1, 1, 2048],
+ [3, 3],
+ [2, 2],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ False,
+ ),
+ (
+ [1, 1, 1, 2048],
+ [4, 4],
+ [2, 2],
+ [2, 3],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ False,
+ ),
+ # Use 'te' schedule
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ True,
+ ),
+ (
+ [1, 1, 1, 2048],
+ [8, 8],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ True,
+ ),
+ )
+
+ @tvm.testing.fixture
+ def expected_output_np(
+ self,
+ input_np,
+ kernel,
+ stride,
+ dilation,
+ padding,
+ ceil_mode,
+ count_include_pad,
+ ):
+ pad_before = padding[:2]
+ pad_after = padding[2:]
+ ref_np = tvm.topi.testing.poolnd_python(
+ input_np,
+ kernel,
+ stride,
+ dilation,
+ pad_before,
+ pad_after,
+ "avg", # pool_type
+ count_include_pad,
+ False, # ceil_mode,
+ layout="NHWC",
+ )
+ return ref_np
+
+ @tvm.testing.fixture
+ def input_shape(self, output_shape, kernel, padding, stride, dilation, output_layout):
+ # Input shape without crouton padding; 'ceil' is being ignored from calculation:
+ o_b, o_h, o_w, o_c = output_shape
+ d_h, d_w = dilation
+ s_h, s_w = stride
+ k_h, k_w = kernel
+ pad_before_h, pad_before_w = padding[:2]
+ pad_after_h, pad_after_w = padding[2:]
+
+ if output_layout == "n11c-1024c":
+ assert (
+ pad_before_w == 0 and pad_after_w == 0 and pad_before_h == 0 and pad_after_h == 0
+ ), "Padding must be zero for n11c-1024c layout!!"
Review Comment:
Nit: (at most) one exclamation point is probably enough here. Having this appear in an assertion-failed message is probably enough to get the user's attention.
--
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] cconvey commented on a diff in pull request #11417: Implement avg_pool2d slice op
Posted by GitBox <gi...@apache.org>.
cconvey commented on code in PR #11417:
URL: https://github.com/apache/tvm/pull/11417#discussion_r882086669
##########
python/tvm/topi/hexagon/utils.py:
##########
@@ -0,0 +1,75 @@
+# 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.
+
+from tvm import te
+
+
+def n11c_1024c_2d(n, h, w, c):
+ return [n, h, w, c // 1024, te.AXIS_SEPARATOR, c % 1024]
+
+
+def n11c_1024c_1d(n, h, w, c):
+ return [n, h, w, c // 1024, c % 1024]
+
+
+def nhwc_8h2w32c2w_2d(n, h, w, c):
+ return [n, h // 8, w // 4, c // 32, te.AXIS_SEPARATOR, h % 8, (w % 4) // 2, c % 32, w % 2]
+
+
+def nhwc_8h2w32c2w_1d(n, h, w, c):
+ return [n, h // 8, w // 4, c // 32, h % 8, (w % 4) // 2, c % 32, w % 2]
Review Comment:
Would these functions' purpose be clearer if their names started with something like `get_shape_...`, `xform_layout_...`, etc?
--
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] cconvey commented on a diff in pull request #11417: Implement avg_pool2d slice op
Posted by GitBox <gi...@apache.org>.
cconvey commented on code in PR #11417:
URL: https://github.com/apache/tvm/pull/11417#discussion_r882086669
##########
python/tvm/topi/hexagon/utils.py:
##########
@@ -0,0 +1,75 @@
+# 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.
+
+from tvm import te
+
+
+def n11c_1024c_2d(n, h, w, c):
+ return [n, h, w, c // 1024, te.AXIS_SEPARATOR, c % 1024]
+
+
+def n11c_1024c_1d(n, h, w, c):
+ return [n, h, w, c // 1024, c % 1024]
+
+
+def nhwc_8h2w32c2w_2d(n, h, w, c):
+ return [n, h // 8, w // 4, c // 32, te.AXIS_SEPARATOR, h % 8, (w % 4) // 2, c % 32, w % 2]
+
+
+def nhwc_8h2w32c2w_1d(n, h, w, c):
+ return [n, h // 8, w // 4, c // 32, h % 8, (w % 4) // 2, c % 32, w % 2]
Review Comment:
Would these functions' purpose be clearer if their names started with something like `get_shape_...`?
--
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] jverma-quic commented on a diff in pull request #11417: [Hexagon] Implement avg_pool2d slice op
Posted by GitBox <gi...@apache.org>.
jverma-quic commented on code in PR #11417:
URL: https://github.com/apache/tvm/pull/11417#discussion_r883953087
##########
tests/python/contrib/test_hexagon/test_avg_pool2d_slice.py:
##########
@@ -0,0 +1,434 @@
+# 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
+import numpy as np
+
+np.set_printoptions(threshold=np.inf)
+from tvm import te, topi
+from tvm.tir.stmt_functor import post_order_visit
+
+import tvm.testing
+from tvm.topi import testing
+from tvm.contrib.hexagon.build import HexagonLauncher
+import tvm.topi.hexagon.slice_ops as sl
+from .infrastructure import allocate_hexagon_array, transform_numpy
+
+
+input_layout = tvm.testing.parameter(
+ "nhwc-8h2w32c2w",
+)
+
+
+@tvm.testing.fixture
+def input_np(input_shape, dtype):
+ return np.random.random(input_shape).astype(dtype)
+
+
+@tvm.testing.fixture
+def transformed_expected_output_np(expected_output_np, output_layout):
+ return transform_numpy(expected_output_np, output_layout)
+
+
+@tvm.testing.fixture
+def transformed_input_np_padded(input_np_padded, input_layout):
+ return transform_numpy(input_np_padded, input_layout)
+
+
+class TestAvgPool2dSlice:
+ # NOTE: input_layout is always assumed to be "nhwc-8h2w32c2w"
+ (
+ output_shape,
+ kernel,
+ stride,
+ dilation,
+ padding,
+ ceil_mode,
+ count_include_pad,
+ output_layout,
+ dtype,
+ use_te_sched,
+ ) = tvm.testing.parameters(
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 16, 16, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [8, 8],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [1, 1],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ # Test non-one stride and dilation
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [2, 3],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [2, 2],
+ [2, 2],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [2, 2],
+ [2, 3],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ # Test non-zero padding
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [1, 1, 1, 1],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [1, 2, 3, 4],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [1, 2, 3, 4],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [3, 2],
+ [2, 3],
+ [1, 2, 3, 4],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ False,
+ ),
+ # Test n11c-1024c layout which will require input and output to have different layout
+ (
+ [1, 1, 1, 2048],
+ [8, 8],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ False,
+ ),
+ (
+ [1, 1, 1, 2048],
+ [6, 6],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ False,
+ ),
+ (
+ [1, 1, 1, 2048],
+ [3, 3],
+ [2, 2],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ False,
+ ),
+ (
+ [1, 1, 1, 2048],
+ [4, 4],
+ [2, 2],
+ [2, 3],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ False,
+ ),
+ # Use 'te' schedule
+ (
+ [1, 8, 8, 32],
+ [3, 3],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "nhwc-8h2w32c2w",
+ "float16",
+ True,
+ ),
+ (
+ [1, 1, 1, 2048],
+ [8, 8],
+ [1, 1],
+ [1, 1],
+ [0, 0, 0, 0],
+ False,
+ True,
+ "n11c-1024c",
+ "float16",
+ True,
+ ),
+ )
+
+ @tvm.testing.fixture
+ def expected_output_np(
+ self,
+ input_np,
+ kernel,
+ stride,
+ dilation,
+ padding,
+ ceil_mode,
+ count_include_pad,
+ ):
+ pad_before = padding[:2]
+ pad_after = padding[2:]
+ ref_np = tvm.topi.testing.poolnd_python(
+ input_np,
+ kernel,
+ stride,
+ dilation,
+ pad_before,
+ pad_after,
+ "avg", # pool_type
+ count_include_pad,
+ False, # ceil_mode,
+ layout="NHWC",
+ )
+ return ref_np
+
+ @tvm.testing.fixture
+ def input_shape(self, output_shape, kernel, padding, stride, dilation, output_layout):
+ # Input shape without crouton padding; 'ceil' is being ignored from calculation:
+ o_b, o_h, o_w, o_c = output_shape
+ d_h, d_w = dilation
+ s_h, s_w = stride
+ k_h, k_w = kernel
+ pad_before_h, pad_before_w = padding[:2]
+ pad_after_h, pad_after_w = padding[2:]
+
+ if output_layout == "n11c-1024c":
+ assert (
+ pad_before_w == 0 and pad_after_w == 0 and pad_before_h == 0 and pad_after_h == 0
+ ), "Padding must be zero for n11c-1024c layout!!"
+ assert o_h == 1 and o_w == 1, "Output height and width must be 1!"
+
+ in_h = (o_h - 1) * s_h + d_h * (k_h - 1) + 1 - pad_before_h - pad_after_h
+ in_w = (o_w - 1) * s_w + d_w * (k_w - 1) + 1 - pad_before_w - pad_after_w
+
+ return [o_b, in_h, in_w, o_c]
+
+ @tvm.testing.fixture
+ def input_shape_padded(self, input_shape, padding, output_layout):
+ # Input shape with regular and crouton padding.
+ # Input width and height are padded to a multiple of croutons.
+ # NOTE: Input layout is always assumed to be nhwc-8h2w32c2w. Only the output layout can be
+ # different.
+ pad_before_h, pad_before_w = padding[:2]
+ pad_after_h, pad_after_w = padding[2:]
+ padded_input_height = ((input_shape[1] + pad_before_h + pad_after_h + 7) // 8) * 8
+ padded_input_width = ((input_shape[2] + pad_before_w + pad_after_w + 3) // 4) * 4
+ return [input_shape[0], padded_input_height, padded_input_width, input_shape[3]]
+
+ @tvm.testing.fixture
+ def input_np_padded(self, input_np, input_shape, input_shape_padded, padding):
+ pad_before_h, pad_before_w = padding[:2]
+ pad_after_h = (
+ input_shape_padded[1] - input_shape[1] - pad_before_h
+ ) # pad_after for height with crouton padding
+ pad_after_w = (
+ input_shape_padded[2] - input_shape[2] - pad_before_w
+ ) # pad_after for width with crouton padding
+ input_padded = np.pad(
+ input_np,
+ ((0, 0), (pad_before_h, pad_after_h), (pad_before_w, pad_after_w), (0, 0)),
+ "constant",
+ )
+ return input_padded
+
+ @tvm.testing.requires_hexagon
+ def test_avg_pool2d_slice(
+ self,
+ use_te_sched,
+ stride,
+ kernel,
+ dtype,
+ dilation,
+ padding,
+ count_include_pad,
+ input_layout,
+ output_layout,
+ output_shape,
+ input_shape,
+ input_shape_padded,
+ input_np,
+ input_np_padded,
+ transformed_input_np_padded,
+ transformed_expected_output_np,
+ expected_output_np,
+ hexagon_session,
+ ):
+
+ target_hexagon = tvm.target.hexagon("v69")
Review Comment:
I think it should be okay but will confirm at our end. BTW, do you know which Hexagon SDK is used for upstream CI?
--
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