Lunderberg commented on code in PR #11522: URL: https://github.com/apache/tvm/pull/11522#discussion_r887990956
########## tests/python/contrib/test_hexagon/test_batch_flatten.py: ########## @@ -0,0 +1,130 @@ +# 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 numpy as np +import pytest + +import tvm +import tvm.testing +import tvm.topi.hexagon.slice_ops as sl +from tvm import te, topi +from tvm.contrib.hexagon.build import HexagonLauncher +from tvm.topi import testing + +from .infrastructure import allocate_hexagon_array + + +def n11c_1024c_1d(n, h, w, c): + return [n, h, w, c // 1024, tvm.te.AXIS_SEPARATOR, c % 1024] + + +def nc_1024_1d(n, c): + return [n, c // 1024, tvm.te.AXIS_SEPARATOR, c % 1024] + + +def transform_numpy(arr_np, layout): + if layout == "nhwc": + return arr_np + elif layout == "n11c-1024c-1d": + N, H, W, C = arr_np.shape + return arr_np.reshape([N, H, W, C // 1024, 1024]) + elif layout == "nc-1d": + N, C = arr_np.shape + return arr_np.reshape([N, C // 1024, 1024]) + + [email protected] +def transformed_expected_output_np(expected_output_np, output_layout): + return transform_numpy(expected_output_np, output_layout) + + +class BaseTestBatchFlatten: + ( + input_shape, + input_layout, + output_layout, + input_axis_sep, + output_axis_sep, + ) = tvm.testing.parameters( + ((1, 1, 1, 2048), "n11c-1024c-1d", "nc-1d", [4], [2]), + ((1, 2, 4, 2048), "n11c-1024c-1d", "nc-1d", [4], [2]), + ((1, 8, 8, 1024), "n11c-1024c-1d", "nc-1d", [4], [2]), + ((2, 4, 8, 1024), "n11c-1024c-1d", "nc-1d", [4], [2]), + ((2, 3, 5, 2048), "n11c-1024c-1d", "nc-1d", [4], [2]), + ) + data_type = tvm.testing.parameter("float16") + + +class TestBatchFlatten(BaseTestBatchFlatten): + @tvm.testing.fixture + def output_shape(self, input_shape): + return input_shape[0], input_shape[1] * input_shape[2] * input_shape[3] + + @tvm.testing.requires_hexagon + def test_batch_flatten( + self, + data_type, + input_shape, + input_layout, + input_axis_sep, + output_shape, + output_layout, + output_axis_sep, + hexagon_session, + ): + target_hexagon = tvm.target.hexagon("v69") + target = tvm.target.Target(target_hexagon, host=target_hexagon) + A = te.placeholder(input_shape, name="A", dtype=data_type) + D = sl.batch_flatten_compute(A) + tir_s = sl.batch_flatten_stir_schedule( + D, + A, + nc_1024_1d, + n11c_1024c_1d, + ) + func_name = "batch_flatten" + with tvm.transform.PassContext(opt_level=3, config={"tir.disable_assert": True}): + tir_irm = tvm.lower(tir_s.mod, [A, D], name=func_name) + runtime_module = tvm.build(tir_irm, [A, D], target=target, name=func_name) Review Comment: `tvm.build` doesn't require the input to be lowered, as the first step it performs is to call `tvm.lower`. Lowering shouldn't have any effect the second time around, but isn't required. ########## tests/python/contrib/test_hexagon/test_batch_flatten.py: ########## @@ -0,0 +1,130 @@ +# 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 numpy as np +import pytest + +import tvm +import tvm.testing +import tvm.topi.hexagon.slice_ops as sl +from tvm import te, topi +from tvm.contrib.hexagon.build import HexagonLauncher +from tvm.topi import testing + +from .infrastructure import allocate_hexagon_array + + +def n11c_1024c_1d(n, h, w, c): + return [n, h, w, c // 1024, tvm.te.AXIS_SEPARATOR, c % 1024] + + +def nc_1024_1d(n, c): + return [n, c // 1024, tvm.te.AXIS_SEPARATOR, c % 1024] + + +def transform_numpy(arr_np, layout): + if layout == "nhwc": + return arr_np + elif layout == "n11c-1024c-1d": + N, H, W, C = arr_np.shape + return arr_np.reshape([N, H, W, C // 1024, 1024]) + elif layout == "nc-1d": + N, C = arr_np.shape + return arr_np.reshape([N, C // 1024, 1024]) + + [email protected] +def transformed_expected_output_np(expected_output_np, output_layout): + return transform_numpy(expected_output_np, output_layout) + + +class BaseTestBatchFlatten: + ( + input_shape, + input_layout, + output_layout, + input_axis_sep, + output_axis_sep, + ) = tvm.testing.parameters( + ((1, 1, 1, 2048), "n11c-1024c-1d", "nc-1d", [4], [2]), + ((1, 2, 4, 2048), "n11c-1024c-1d", "nc-1d", [4], [2]), + ((1, 8, 8, 1024), "n11c-1024c-1d", "nc-1d", [4], [2]), + ((2, 4, 8, 1024), "n11c-1024c-1d", "nc-1d", [4], [2]), + ((2, 3, 5, 2048), "n11c-1024c-1d", "nc-1d", [4], [2]), + ) + data_type = tvm.testing.parameter("float16") + + +class TestBatchFlatten(BaseTestBatchFlatten): + @tvm.testing.fixture + def output_shape(self, input_shape): + return input_shape[0], input_shape[1] * input_shape[2] * input_shape[3] + + @tvm.testing.requires_hexagon + def test_batch_flatten( + self, + data_type, + input_shape, + input_layout, + input_axis_sep, + output_shape, + output_layout, + output_axis_sep, + hexagon_session, + ): + target_hexagon = tvm.target.hexagon("v69") + target = tvm.target.Target(target_hexagon, host=target_hexagon) + A = te.placeholder(input_shape, name="A", dtype=data_type) + D = sl.batch_flatten_compute(A) + tir_s = sl.batch_flatten_stir_schedule( + D, + A, + nc_1024_1d, + n11c_1024c_1d, + ) + func_name = "batch_flatten" + with tvm.transform.PassContext(opt_level=3, config={"tir.disable_assert": True}): + tir_irm = tvm.lower(tir_s.mod, [A, D], name=func_name) Review Comment: The argument list passed to `tvm.lower` and `tvm.build` (here, `[A, D]`) is used to specify which TE tensors should be exposed as TIR arguments, and should only be provided for TE-based schedules. For STIR, the arguments are already defined in the call to `te.create_prim_func`. ########## 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) Review Comment: Good catch, and thank you for finding this. I missed passing this through in #10904. ########## tests/python/contrib/test_hexagon/test_batch_flatten.py: ########## @@ -0,0 +1,130 @@ +# 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 numpy as np +import pytest + +import tvm +import tvm.testing +import tvm.topi.hexagon.slice_ops as sl +from tvm import te, topi +from tvm.contrib.hexagon.build import HexagonLauncher +from tvm.topi import testing + +from .infrastructure import allocate_hexagon_array + + +def n11c_1024c_1d(n, h, w, c): + return [n, h, w, c // 1024, tvm.te.AXIS_SEPARATOR, c % 1024] + + +def nc_1024_1d(n, c): + return [n, c // 1024, tvm.te.AXIS_SEPARATOR, c % 1024] + + +def transform_numpy(arr_np, layout): + if layout == "nhwc": + return arr_np + elif layout == "n11c-1024c-1d": + N, H, W, C = arr_np.shape + return arr_np.reshape([N, H, W, C // 1024, 1024]) + elif layout == "nc-1d": + N, C = arr_np.shape + return arr_np.reshape([N, C // 1024, 1024]) + + [email protected] +def transformed_expected_output_np(expected_output_np, output_layout): + return transform_numpy(expected_output_np, output_layout) + + +class BaseTestBatchFlatten: Review Comment: Separating parameters out into a base class for testing isn't required. I use this if I want to define several groups of parameters with associated names (e.g. input sizes used by a specific model), but since there's only one subclass of `BaseTestBatchFlatten`, I'd move the contents into `TestBatchFlatten` and remove `BaseTestBatchFlatten` altogether. There's also a lot of repetition between the test cases, so it isn't immediately clear that these are only varying the input shape and not the layouts. I'd rearrange them so that the shape is emphasized as the parameter being varied. ```python input_shape = tvm.testing.parameter( (1, 1, 1, 2048), (1, 2, 4, 2048), (1, 8, 8, 1024), (2, 4, 8, 1024), (2, 3, 5, 2048), ) input_layout, input_axis_sep = tvm.testing.parameters(("n11c-1024c-1d", [4])) output_layout, output_axis_sep = tvm.testing.parameters( ("nc-1d", [2])) data_type = tvm.testing.parameter("float16") ``` -- This is an automated message from the Apache Git Service. To respond to the message, please log on to GitHub and use the URL above to go to the specific comment. To unsubscribe, e-mail: [email protected] For queries about this service, please contact Infrastructure at: [email protected]
