From a8c4e8712ee4f7308a80e7358dfce4ee6b67553a Mon Sep 17 00:00:00 2001 From: "Coplin, Jared" Date: Thu, 2 Jun 2022 12:51:46 -0500 Subject: [PATCH 01/13] [HEXAGON] Initial clip operator for Hexagon --- python/tvm/topi/hexagon/slice_ops/__init__.py | 22 +++ python/tvm/topi/hexagon/slice_ops/clip.py | 62 +++++++++ python/tvm/topi/hexagon/utils.py | 75 ++++++++++ .../contrib/test_hexagon/infrastructure.py | 16 ++- .../python/contrib/test_hexagon/test_clip.py | 128 ++++++++++++++++++ 5 files changed, 302 insertions(+), 1 deletion(-) create mode 100755 python/tvm/topi/hexagon/slice_ops/__init__.py create mode 100755 python/tvm/topi/hexagon/slice_ops/clip.py create mode 100644 python/tvm/topi/hexagon/utils.py create mode 100755 tests/python/contrib/test_hexagon/test_clip.py diff --git a/python/tvm/topi/hexagon/slice_ops/__init__.py b/python/tvm/topi/hexagon/slice_ops/__init__.py new file mode 100755 index 000000000000..80598fcdf3db --- /dev/null +++ b/python/tvm/topi/hexagon/slice_ops/__init__.py @@ -0,0 +1,22 @@ +# 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. + +""" Computes and Schedules for Hexagon slice ops. """ + +# pylint: disable=wildcard-import + +from .clip import * diff --git a/python/tvm/topi/hexagon/slice_ops/clip.py b/python/tvm/topi/hexagon/slice_ops/clip.py new file mode 100755 index 000000000000..8604e8e59058 --- /dev/null +++ b/python/tvm/topi/hexagon/slice_ops/clip.py @@ -0,0 +1,62 @@ +# 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 + +""" +Clip the elements in `A` between `A_min` and `A_max`. +""" + +from tvm.ir.module import IRModule +from tvm import te, tir, topi +from ..utils import get_layout_transform_fn + + +def clip_compute(A, A_min, A_max): + # Use topi implementation + return topi.clip(A, A_min, A_max) + + +def clip_schedule(outs, ins, output_layout: str, input_layout: str): + A = ins + M = outs + + func = te.create_prim_func([A, M]) + + s = tir.Schedule(func) + + block = s.get_block("compute") + + input_transformed_layout = get_layout_transform_fn(input_layout) + s.transform_layout(block, buffer=("read", 0), index_map=input_transformed_layout) + + output_transformed_layout = get_layout_transform_fn(output_layout) + s.transform_layout(block, buffer=("write", 0), index_map=output_transformed_layout) + + n, h, w, c = s.get_loops(block) + + ho, hi = s.split(h, [None, 8]) + wo, wi = s.split(w, [None, 4]) + co, ci = s.split(c, [None, 32]) + wio, wii = s.split(wi, [None, 2]) + + s.reorder(n, ho, wo, co, hi, wio, ci, wii) + + fused = s.fuse(ci, wii) + s.vectorize(fused) + + return s diff --git a/python/tvm/topi/hexagon/utils.py b/python/tvm/topi/hexagon/utils.py new file mode 100644 index 000000000000..1dc84139ae70 --- /dev/null +++ b/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) + if axis_separators: + s.set_axis_separator(block, block_index, buffer_type, axis_separators) diff --git a/tests/python/contrib/test_hexagon/infrastructure.py b/tests/python/contrib/test_hexagon/infrastructure.py index 0c9a9478c870..b9270df7ef0e 100644 --- a/tests/python/contrib/test_hexagon/infrastructure.py +++ b/tests/python/contrib/test_hexagon/infrastructure.py @@ -48,7 +48,7 @@ def allocate_hexagon_array( for dim_i, dim_f in zip(boundaries[:-1], boundaries[1:]) ] - arr = tvm.nd.empty(physical_shape, dtype=dtype, device=dev) + arr = tvm.nd.empty(physical_shape, dtype=dtype, device=dev, mem_scope=mem_scope) if data is not None: arr.copyfrom(data.reshape(physical_shape)) @@ -228,3 +228,17 @@ 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) + else: + raise RuntimeError(f"Unexpected layout '{layout}'") diff --git a/tests/python/contrib/test_hexagon/test_clip.py b/tests/python/contrib/test_hexagon/test_clip.py new file mode 100755 index 000000000000..4cd5fa41f50e --- /dev/null +++ b/tests/python/contrib/test_hexagon/test_clip.py @@ -0,0 +1,128 @@ +# 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 + +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): + # transform_numpy does not support layouts that end in -*d + return transform_numpy(expected_output_np, output_layout[:-3]) + + +@tvm.testing.fixture +def transformed_input_np(input_np, input_layout): + # transform_numpy does not support layouts that end in -*d + return transform_numpy(input_np, input_layout[:-3]) + + +class TestClipSlice: + input_shape, output_shape, A_min, A_max, output_layout, dtype = tvm.testing.parameters( + ([1, 256, 256, 32], [1, 256, 256, 32], 0.1, 0.5, "nhwc-8h2w32c2w-2d", "float16") + ) + + @tvm.testing.fixture + def expected_output_np(self, input_np, A_min, A_max): + ref_np = np.clip(input_np, A_min, A_max) + return ref_np + + def test_clip_slice( + self, + input_shape, + output_shape, + input_np, + input_layout, + output_layout, + dtype, + A_min, + A_max, + transformed_input_np, + transformed_expected_output_np, + hexagon_session, + ): + # establish target and input placeholder + target_hexagon = tvm.target.hexagon("v68") + A = te.placeholder(input_shape, name="A", dtype=dtype) + + # get the compute function and schedule + M = sl.clip_compute(A, A_min, A_max) + + # Assume layout is nhwc-8h2w32c2w-2d + tir_schedule = sl.clip_schedule(M, A, output_layout, input_layout) + sch = tir_schedule.mod + + # build the function + with tvm.transform.PassContext(opt_level=3): + func = tvm.build( + sch, [A, M], tvm.target.Target(target_hexagon, host=target_hexagon), name="clip" + ) + + # allocate input and output nd arrays + axis_separators = [4] + input_arr = allocate_hexagon_array( + hexagon_session.device, + data=transformed_input_np, + dtype=dtype, + axis_separators=axis_separators, + mem_scope="global.vtcm", + ) + + output_arr = allocate_hexagon_array( + hexagon_session.device, + transformed_expected_output_np.shape, + dtype=dtype, + axis_separators=axis_separators, + mem_scope="global.vtcm", + ) + + # execute + mod = hexagon_session.load_module(func) + mod(input_arr, output_arr) + + # convert output nd array to numpy array + output_np = output_arr.numpy() + b, h, w, c = output_shape + reshaped_output_np = np.reshape(output_np, [b, h // 8, w // 4, c // 32, 8, 2, 32, 2]) + + # test results + np.testing.assert_allclose( + reshaped_output_np, transformed_expected_output_np, rtol=1e-3, atol=1e-3 + ) + + +if __name__ == "__main__": + sys.exit(pytest.main(sys.argv)) From 0e4fddcc6c98f5ed5c6d4c9c8c0a662d5e6e85dd Mon Sep 17 00:00:00 2001 From: "Coplin, Jared" Date: Thu, 2 Jun 2022 13:14:58 -0500 Subject: [PATCH 02/13] Changes to utils and infra for pylint --- python/tvm/topi/hexagon/utils.py | 43 +++++-------------- .../contrib/test_hexagon/infrastructure.py | 33 ++++++++------ 2 files changed, 29 insertions(+), 47 deletions(-) diff --git a/python/tvm/topi/hexagon/utils.py b/python/tvm/topi/hexagon/utils.py index 1dc84139ae70..af6e3de9c350 100644 --- a/python/tvm/topi/hexagon/utils.py +++ b/python/tvm/topi/hexagon/utils.py @@ -14,62 +14,39 @@ # KIND, either express or implied. See the License for the # specific language governing permissions and limitations # under the License. - +# pylint: disable=invalid-name +"""Common hexagon specific utilities""" from tvm import te def n11c_1024c_2d(n, h, w, c): + """Return index map for n11c_1024 2d layout""" return [n, h, w, c // 1024, te.AXIS_SEPARATOR, c % 1024] def n11c_1024c_1d(n, h, w, c): + """Return index map for n11c_1024 1d layout""" return [n, h, w, c // 1024, c % 1024] def nhwc_8h2w32c2w_2d(n, h, w, c): + """Return index map for nhwc_8h2w32c2w 2d layout""" 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 index map for nhwc_8h2w32c2w 1d layout""" return [n, h // 8, w // 4, c // 32, h % 8, (w % 4) // 2, c % 32, w % 2] def get_layout_transform_fn(layout): + """Return index map function as per the layout string""" if layout == "nhwc-8h2w32c2w-2d": return nhwc_8h2w32c2w_2d if layout == "nhwc-8h2w32c2w-1d": return nhwc_8h2w32c2w_1d - elif layout == "n11c-1024c-2d": + if layout == "n11c-1024c-2d": return n11c_1024c_2d - elif layout == "n11c-1024c-1d": + if 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) - if axis_separators: - s.set_axis_separator(block, block_index, buffer_type, axis_separators) + raise RuntimeError(f"Unexpected layout '{layout}'") diff --git a/tests/python/contrib/test_hexagon/infrastructure.py b/tests/python/contrib/test_hexagon/infrastructure.py index b9270df7ef0e..a6e1fa662e4a 100644 --- a/tests/python/contrib/test_hexagon/infrastructure.py +++ b/tests/python/contrib/test_hexagon/infrastructure.py @@ -14,6 +14,7 @@ # KIND, either express or implied. See the License for the # specific language governing permissions and limitations # under the License. +# pylint: disable=invalid-name """ Hexagon testing infrastructure """ @@ -47,7 +48,6 @@ def allocate_hexagon_array( numpy.prod(tensor_shape[dim_i:dim_f]) for dim_i, dim_f in zip(boundaries[:-1], boundaries[1:]) ] - arr = tvm.nd.empty(physical_shape, dtype=dtype, device=dev, mem_scope=mem_scope) if data is not None: @@ -229,16 +229,21 @@ 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) - else: - raise RuntimeError(f"Unexpected layout '{layout}'") + +def transform_numpy(arr_np, current_layout: str, new_layout: str): + """Reshape and transpose numpy array according to the specified layout""" + if current_layout == "nhwc": + if new_layout == "nhwc": + return arr_np + if new_layout in ["nhwc-8h2w32c2w-2d", "nhwc-8h2w32c2w-1d"]: + 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 + ) + if new_layout in ["n11c-1024c-2d", "n11c-1024c-1d"]: + 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, 1, 1, c // 1024, 1024]) + + raise RuntimeError(f"Unexpected new_layout '{new_layout}'") + raise RuntimeError(f"Unexpected current_layout '{current_layout}'") From 9fa366826a69f5ae8412424db8e32d6e33470dab Mon Sep 17 00:00:00 2001 From: "Coplin, Jared" Date: Thu, 2 Jun 2022 15:02:03 -0500 Subject: [PATCH 03/13] Remove unused import --- python/tvm/topi/hexagon/slice_ops/clip.py | 1 - 1 file changed, 1 deletion(-) diff --git a/python/tvm/topi/hexagon/slice_ops/clip.py b/python/tvm/topi/hexagon/slice_ops/clip.py index 8604e8e59058..3c64e2847e33 100755 --- a/python/tvm/topi/hexagon/slice_ops/clip.py +++ b/python/tvm/topi/hexagon/slice_ops/clip.py @@ -21,7 +21,6 @@ Clip the elements in `A` between `A_min` and `A_max`. """ -from tvm.ir.module import IRModule from tvm import te, tir, topi from ..utils import get_layout_transform_fn From adf2e47b590cd118b283cd32cbb347939a37edea Mon Sep 17 00:00:00 2001 From: "Coplin, Jared" Date: Thu, 2 Jun 2022 15:03:06 -0500 Subject: [PATCH 04/13] Use tvm.testing.main() --- tests/python/contrib/test_hexagon/test_clip.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/python/contrib/test_hexagon/test_clip.py b/tests/python/contrib/test_hexagon/test_clip.py index 4cd5fa41f50e..9b3dd0bfd2ad 100755 --- a/tests/python/contrib/test_hexagon/test_clip.py +++ b/tests/python/contrib/test_hexagon/test_clip.py @@ -125,4 +125,4 @@ def test_clip_slice( if __name__ == "__main__": - sys.exit(pytest.main(sys.argv)) + tvm.testing.main() From 7b3498398844f91b63db06c5a3c413ce26fb858d Mon Sep 17 00:00:00 2001 From: "Coplin, Jared" Date: Thu, 2 Jun 2022 15:12:48 -0500 Subject: [PATCH 05/13] Address pylint error --- python/tvm/topi/hexagon/slice_ops/clip.py | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/python/tvm/topi/hexagon/slice_ops/clip.py b/python/tvm/topi/hexagon/slice_ops/clip.py index 3c64e2847e33..2beb2df643bb 100755 --- a/python/tvm/topi/hexagon/slice_ops/clip.py +++ b/python/tvm/topi/hexagon/slice_ops/clip.py @@ -26,11 +26,16 @@ def clip_compute(A, A_min, A_max): - # Use topi implementation + """ + Use topi clip implementation + """ return topi.clip(A, A_min, A_max) def clip_schedule(outs, ins, output_layout: str, input_layout: str): + """ + Hexagon clip schedule + """ A = ins M = outs From ff1279fa171da5d3267891fa1a013c40a06780c4 Mon Sep 17 00:00:00 2001 From: "Coplin, Jared" Date: Fri, 3 Jun 2022 12:13:23 -0500 Subject: [PATCH 06/13] Fix incorrect function call --- tests/python/contrib/test_hexagon/test_clip.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tests/python/contrib/test_hexagon/test_clip.py b/tests/python/contrib/test_hexagon/test_clip.py index 9b3dd0bfd2ad..92e27bf10785 100755 --- a/tests/python/contrib/test_hexagon/test_clip.py +++ b/tests/python/contrib/test_hexagon/test_clip.py @@ -41,13 +41,13 @@ def input_np(input_shape, dtype): @tvm.testing.fixture def transformed_expected_output_np(expected_output_np, output_layout): # transform_numpy does not support layouts that end in -*d - return transform_numpy(expected_output_np, output_layout[:-3]) + return transform_numpy(expected_output_np, "nhwc", output_layout[:-3]) @tvm.testing.fixture def transformed_input_np(input_np, input_layout): # transform_numpy does not support layouts that end in -*d - return transform_numpy(input_np, input_layout[:-3]) + return transform_numpy(input_np, "nhwc", input_layout[:-3]) class TestClipSlice: From b38b0e974cc80a28fe69fc70794ffd98658d4a4e Mon Sep 17 00:00:00 2001 From: "Coplin, Jared" Date: Fri, 3 Jun 2022 14:56:24 -0500 Subject: [PATCH 07/13] Changes to calls to transform_numpy --- tests/python/contrib/test_hexagon/test_clip.py | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/tests/python/contrib/test_hexagon/test_clip.py b/tests/python/contrib/test_hexagon/test_clip.py index 92e27bf10785..4a33c98f36c2 100755 --- a/tests/python/contrib/test_hexagon/test_clip.py +++ b/tests/python/contrib/test_hexagon/test_clip.py @@ -40,14 +40,12 @@ def input_np(input_shape, dtype): @tvm.testing.fixture def transformed_expected_output_np(expected_output_np, output_layout): - # transform_numpy does not support layouts that end in -*d - return transform_numpy(expected_output_np, "nhwc", output_layout[:-3]) + return transform_numpy(expected_output_np, "nhwc", output_layout) @tvm.testing.fixture def transformed_input_np(input_np, input_layout): - # transform_numpy does not support layouts that end in -*d - return transform_numpy(input_np, "nhwc", input_layout[:-3]) + return transform_numpy(input_np, "nhwc", input_layout) class TestClipSlice: From 8ac841af1ff99d3db075925e278bca77c7deab01 Mon Sep 17 00:00:00 2001 From: "Coplin, Jared" Date: Thu, 16 Jun 2022 10:48:06 -0500 Subject: [PATCH 08/13] Add newline at end of file --- python/tvm/topi/hexagon/slice_ops/__init__.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/tvm/topi/hexagon/slice_ops/__init__.py b/python/tvm/topi/hexagon/slice_ops/__init__.py index c64ae4abe11e..14d7e717747c 100755 --- a/python/tvm/topi/hexagon/slice_ops/__init__.py +++ b/python/tvm/topi/hexagon/slice_ops/__init__.py @@ -20,4 +20,4 @@ # pylint: disable=wildcard-import from .avg_pool2d import avg_pool2d_compute, avg_pool2d_STIR_schedule -from .clip import * \ No newline at end of file +from .clip import * From 5657b4e651dbd2a87b82ef603f56763b2951d939 Mon Sep 17 00:00:00 2001 From: "Coplin, Jared" Date: Tue, 21 Jun 2022 15:11:32 -0500 Subject: [PATCH 09/13] Add requires_hexagon and rename under topi --- tests/python/contrib/test_hexagon/{ => topi}/test_clip.py | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) rename tests/python/contrib/test_hexagon/{ => topi}/test_clip.py (97%) diff --git a/tests/python/contrib/test_hexagon/test_clip.py b/tests/python/contrib/test_hexagon/topi/test_clip.py similarity index 97% rename from tests/python/contrib/test_hexagon/test_clip.py rename to tests/python/contrib/test_hexagon/topi/test_clip.py index 4a33c98f36c2..c129fc169574 100755 --- a/tests/python/contrib/test_hexagon/test_clip.py +++ b/tests/python/contrib/test_hexagon/topi/test_clip.py @@ -26,7 +26,7 @@ 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 +from ..infrastructure import allocate_hexagon_array, transform_numpy input_layout = tvm.testing.parameter( "nhwc-8h2w32c2w-2d", @@ -58,6 +58,7 @@ def expected_output_np(self, input_np, A_min, A_max): ref_np = np.clip(input_np, A_min, A_max) return ref_np + @tvm.testing.requires_hexagon def test_clip_slice( self, input_shape, From 25d89df5946f831243c181616370266ee57896c9 Mon Sep 17 00:00:00 2001 From: "Coplin, Jared" Date: Wed, 22 Jun 2022 09:49:50 -0500 Subject: [PATCH 10/13] Whitespace fix and reduce input size --- tests/python/contrib/test_hexagon/infrastructure.py | 1 + tests/python/contrib/test_hexagon/topi/test_clip.py | 2 +- 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/tests/python/contrib/test_hexagon/infrastructure.py b/tests/python/contrib/test_hexagon/infrastructure.py index a6e1fa662e4a..57a9dff8b424 100644 --- a/tests/python/contrib/test_hexagon/infrastructure.py +++ b/tests/python/contrib/test_hexagon/infrastructure.py @@ -48,6 +48,7 @@ def allocate_hexagon_array( numpy.prod(tensor_shape[dim_i:dim_f]) for dim_i, dim_f in zip(boundaries[:-1], boundaries[1:]) ] + arr = tvm.nd.empty(physical_shape, dtype=dtype, device=dev, mem_scope=mem_scope) if data is not None: diff --git a/tests/python/contrib/test_hexagon/topi/test_clip.py b/tests/python/contrib/test_hexagon/topi/test_clip.py index c129fc169574..8dc7dcf9223e 100755 --- a/tests/python/contrib/test_hexagon/topi/test_clip.py +++ b/tests/python/contrib/test_hexagon/topi/test_clip.py @@ -50,7 +50,7 @@ def transformed_input_np(input_np, input_layout): class TestClipSlice: input_shape, output_shape, A_min, A_max, output_layout, dtype = tvm.testing.parameters( - ([1, 256, 256, 32], [1, 256, 256, 32], 0.1, 0.5, "nhwc-8h2w32c2w-2d", "float16") + ([1, 8, 4, 32], [1, 8, 4, 32], 0.1, 0.5, "nhwc-8h2w32c2w-2d", "float16") ) @tvm.testing.fixture From f03351f3a96509239cbf56517c90ab59fe05b54e Mon Sep 17 00:00:00 2001 From: "Coplin, Jared" Date: Thu, 23 Jun 2022 09:38:45 -0500 Subject: [PATCH 11/13] Remove te tensor arguments --- tests/python/contrib/test_hexagon/topi/test_clip.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tests/python/contrib/test_hexagon/topi/test_clip.py b/tests/python/contrib/test_hexagon/topi/test_clip.py index 8dc7dcf9223e..2c460c390a32 100755 --- a/tests/python/contrib/test_hexagon/topi/test_clip.py +++ b/tests/python/contrib/test_hexagon/topi/test_clip.py @@ -74,7 +74,7 @@ def test_clip_slice( hexagon_session, ): # establish target and input placeholder - target_hexagon = tvm.target.hexagon("v68") + target_hexagon = tvm.target.hexagon("v69") A = te.placeholder(input_shape, name="A", dtype=dtype) # get the compute function and schedule @@ -87,7 +87,7 @@ def test_clip_slice( # build the function with tvm.transform.PassContext(opt_level=3): func = tvm.build( - sch, [A, M], tvm.target.Target(target_hexagon, host=target_hexagon), name="clip" + sch, tvm.target.Target(target_hexagon, host=target_hexagon), name="clip" ) # allocate input and output nd arrays From 9adaa7d1c34d3b819915451099f8bd24ba38f2e9 Mon Sep 17 00:00:00 2001 From: "Coplin, Jared" Date: Mon, 27 Jun 2022 10:58:53 -0500 Subject: [PATCH 12/13] Correct call to tvm.build --- tests/python/contrib/test_hexagon/topi/test_clip.py | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/tests/python/contrib/test_hexagon/topi/test_clip.py b/tests/python/contrib/test_hexagon/topi/test_clip.py index 2c460c390a32..b44c60e93c0c 100755 --- a/tests/python/contrib/test_hexagon/topi/test_clip.py +++ b/tests/python/contrib/test_hexagon/topi/test_clip.py @@ -82,12 +82,11 @@ def test_clip_slice( # Assume layout is nhwc-8h2w32c2w-2d tir_schedule = sl.clip_schedule(M, A, output_layout, input_layout) - sch = tir_schedule.mod # build the function with tvm.transform.PassContext(opt_level=3): func = tvm.build( - sch, tvm.target.Target(target_hexagon, host=target_hexagon), name="clip" + tir_schedule.mod, target=tvm.target.Target(target_hexagon, host=target_hexagon), name="clip" ) # allocate input and output nd arrays From 15cbf67be807cc9d127b9360cc2063130aa6951f Mon Sep 17 00:00:00 2001 From: "Coplin, Jared" Date: Mon, 27 Jun 2022 11:28:24 -0500 Subject: [PATCH 13/13] Run black formatting --- python/tvm/topi/hexagon/slice_ops/__init__.py | 1 - tests/python/contrib/test_hexagon/topi/test_clip.py | 4 +++- 2 files changed, 3 insertions(+), 2 deletions(-) diff --git a/python/tvm/topi/hexagon/slice_ops/__init__.py b/python/tvm/topi/hexagon/slice_ops/__init__.py index 0b0cf8856984..87af3a767c38 100755 --- a/python/tvm/topi/hexagon/slice_ops/__init__.py +++ b/python/tvm/topi/hexagon/slice_ops/__init__.py @@ -21,4 +21,3 @@ from .add_subtract_multiply import * from .softmax_slice import * from .clip import * - diff --git a/tests/python/contrib/test_hexagon/topi/test_clip.py b/tests/python/contrib/test_hexagon/topi/test_clip.py index b44c60e93c0c..37146b55dc1e 100755 --- a/tests/python/contrib/test_hexagon/topi/test_clip.py +++ b/tests/python/contrib/test_hexagon/topi/test_clip.py @@ -86,7 +86,9 @@ def test_clip_slice( # build the function with tvm.transform.PassContext(opt_level=3): func = tvm.build( - tir_schedule.mod, target=tvm.target.Target(target_hexagon, host=target_hexagon), name="clip" + tir_schedule.mod, + target=tvm.target.Target(target_hexagon, host=target_hexagon), + name="clip", ) # allocate input and output nd arrays