From 190e4d93739dd06ff1af7ee9032a7416de70f6af Mon Sep 17 00:00:00 2001 From: Arun Rangasamy Date: Fri, 3 Jun 2022 21:17:55 +0530 Subject: [PATCH 01/11] [Topi][Hexagon] Implement Cast F32ToF16 and F16ToF32 Slice Op --- python/tvm/topi/hexagon/slice_ops/__init__.py | 22 ++ python/tvm/topi/hexagon/slice_ops/cast.py | 188 ++++++++++++++++ .../contrib/test_hexagon/infrastructure.py | 2 +- .../contrib/test_hexagon/test_cast_slice.py | 209 ++++++++++++++++++ 4 files changed, 420 insertions(+), 1 deletion(-) create mode 100644 python/tvm/topi/hexagon/slice_ops/__init__.py create mode 100644 python/tvm/topi/hexagon/slice_ops/cast.py create mode 100644 tests/python/contrib/test_hexagon/test_cast_slice.py diff --git a/python/tvm/topi/hexagon/slice_ops/__init__.py b/python/tvm/topi/hexagon/slice_ops/__init__.py new file mode 100644 index 000000000000..bd55ba289857 --- /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 .cast import * diff --git a/python/tvm/topi/hexagon/slice_ops/cast.py b/python/tvm/topi/hexagon/slice_ops/cast.py new file mode 100644 index 000000000000..3e8024da733a --- /dev/null +++ b/python/tvm/topi/hexagon/slice_ops/cast.py @@ -0,0 +1,188 @@ +# 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. +""" Hexagon slice cast op compute and schedule""" + +from tvm import te +from tvm import tir +from tvm.tir import IndexMap + +# pylint: disable=invalid-name + + +def layout_transform_nhwc_8h2w32c2w(n, h, w, c): + return [ + n, + h // 8, + w // 4, + c // 32, + IndexMap.AXIS_SEPARATOR, + h % 8, + (w % 4) // 2, + c % 32, + w % 2, + ] + + +def layout_transform_nc_1024c(n, c): + return [ + n, + c // 1024, + IndexMap.AXIS_SEPARATOR, + c % 1024, + ] + + +def layout_transform_nhwc_4h2w32c2w(n, h, w, c): + return [ + n, + h // 4, + w // 4, + c // 32, + IndexMap.AXIS_SEPARATOR, + h % 4, + (w % 4) // 2, + (c % 32), + w % 2, + ] + + +def layout_transform_nc_512c(n, c): + return [ + n, + c // 512, + IndexMap.AXIS_SEPARATOR, + c % 512, + ] + + +def get_layout_transform_for_f32(f32_layout_string): + """ + Given f32 layout string, return transform_layout function and + channel/height split factor to be used for scheduling + """ + if f32_layout_string == "nhwc-8h2w32c2w-2d": + return [layout_transform_nhwc_8h2w32c2w, 8] + if f32_layout_string == "nhwc-4h2w32c2w-2d": + return [layout_transform_nhwc_4h2w32c2w, 4] + if f32_layout_string == "nc-1024c-2d": + return [layout_transform_nc_1024c, 1024] + if f32_layout_string == "nc-512c-2d": + return [layout_transform_nc_512c, 512] + raise RuntimeError(f"Unexpected f32_layout '{f32_layout_string}'") + + +def cast_f16_f32_compute(in_tensor): + out_tensor = te.compute( + in_tensor.shape, lambda *indices: in_tensor[indices].astype("float32"), name="CastF16F32" + ) + return out_tensor + + +def cast_f16_f32_stir_schedule_nhwc(func, in_layout, out_layout, h_split_factor): + """Schedule for nhwc f16 to f32 cast: nhwc layout""" + sch = tir.Schedule(func, debug_mask="all") + block_name = "CastF16F32" + n, h, w, c = sch.get_loops(sch.get_block(block_name)) + h_outer, h_inner = sch.split(h, [None, h_split_factor]) + w_outer, w_inner = sch.split(w, [None, 4]) + c_outer, c_inner = sch.split(c, [None, 32]) + w_inner_o, w_inner_i = sch.split(w_inner, [None, 2]) + sch.reorder(n, h_outer, w_outer, c_outer, h_inner, w_inner_o, c_inner, w_inner_i) + sch.transform_layout(block_name, "A", in_layout) + sch.transform_layout(block_name, block_name, out_layout) + fused = sch.fuse(c_inner, w_inner_i) + sch.vectorize(fused) + return sch + + +def cast_f16_f32_stir_schedule_nc(func, in_layout, out_layout, c_split_factor): + """Schedule for nc f16 to f32 cast: nc layout""" + sch = tir.Schedule(func, debug_mask="all") + block_name = "CastF16F32" + _, c = sch.get_loops(sch.get_block(block_name)) + _, c_inner = sch.split(c, [None, c_split_factor]) + sch.transform_layout(block_name, "A", in_layout) + sch.transform_layout(block_name, block_name, out_layout) + sch.vectorize(c_inner) + return sch + + +def cast_f16_f32_schedule(cast_func, in_layout_str, out_layout_str): + """Schedule for f16 to f32 cast: top level function""" + f32_layout_transform_func, split_factor = get_layout_transform_for_f32(out_layout_str) + if in_layout_str == "nhwc-8h2w32c2w-2d": + return cast_f16_f32_stir_schedule_nhwc( + cast_func, + layout_transform_nhwc_8h2w32c2w, + f32_layout_transform_func, + split_factor, + ) + if in_layout_str == "nc-1024c-2d": + return cast_f16_f32_stir_schedule_nc( + cast_func, layout_transform_nc_1024c, f32_layout_transform_func, split_factor + ) + raise RuntimeError(f"Unexpected input_layout, output_layout '{input_layout, output_layout}'") + + +def cast_f32_f16_compute(in_tensor): + out_tensor = te.compute( + in_tensor.shape, lambda *indices: in_tensor[indices].astype("float16"), name="CastF32F16" + ) + return out_tensor + + +def cast_f32_f16_stir_schedule_nhwc(func, in_layout, out_layout, h_split_factor): + """Schedule for nhwc f32 to f16 cast: nhwc layout""" + sch = tir.Schedule(func, debug_mask="all") + block_name = "CastF32F16" + n, h, w, c = sch.get_loops(sch.get_block(block_name)) + h_outer, h_inner = sch.split(h, [None, h_split_factor]) + w_outer, w_inner = sch.split(w, [None, 4]) + c_outer, c_inner = sch.split(c, [None, 32]) + w_inner_o, w_inner_i = sch.split(w_inner, [None, 2]) + sch.reorder(n, h_outer, w_outer, c_outer, h_inner, w_inner_o, c_inner, w_inner_i) + sch.transform_layout(block_name, "A", in_layout) + sch.transform_layout(block_name, block_name, out_layout) + fused = sch.fuse(c_inner, w_inner_i) + sch.vectorize(fused) + return sch + + +def cast_f32_f16_stir_schedule_nc(func, in_layout, out_layout, c_split_factor): + """Schedule for nc f32 to f16 cast: nc layout""" + sch = tir.Schedule(func, debug_mask="all") + block_name = "CastF32F16" + _, c = sch.get_loops(sch.get_block(block_name)) + _, c_inner = sch.split(c, [None, c_split_factor]) + sch.transform_layout(block_name, "A", in_layout) + sch.transform_layout(block_name, block_name, out_layout) + sch.vectorize(c_inner) + return sch + + +def cast_f32_f16_schedule(cast_func, in_layout_str, out_layout_str): + """Schedule for f32 to f16 cast: top level function""" + f32_layout_transform_func, split_factor = get_layout_transform_for_f32(in_layout_str) + if out_layout_str == "nhwc-8h2w32c2w-2d": + return cast_f32_f16_stir_schedule_nhwc( + cast_func, f32_layout_transform_func, layout_transform_nhwc_8h2w32c2w, split_factor + ) + if out_layout_str == "nc-1024c-2d": + return cast_f32_f16_stir_schedule_nc( + cast_func, f32_layout_transform_func, layout_transform_nc_1024c, split_factor + ) + raise RuntimeError(f"Unexpected input_layout, output_layout '{input_layout, output_layout}'") diff --git a/tests/python/contrib/test_hexagon/infrastructure.py b/tests/python/contrib/test_hexagon/infrastructure.py index 0c9a9478c870..01eef86e6b5b 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)) diff --git a/tests/python/contrib/test_hexagon/test_cast_slice.py b/tests/python/contrib/test_hexagon/test_cast_slice.py new file mode 100644 index 000000000000..67cb3ecbde79 --- /dev/null +++ b/tests/python/contrib/test_hexagon/test_cast_slice.py @@ -0,0 +1,209 @@ +# 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. +""" Tests for Hexagon slice cast ops """ +import numpy as np +import pytest + +import tvm +import tvm.testing +from tvm import te, topi +import tvm.topi.hexagon.slice_ops as sl +from tvm.topi import testing +from .infrastructure import allocate_hexagon_array +import tvm.contrib.hexagon + +# pylint: disable=invalid-name + + +def transform_numpy(arr_np, layout): + if layout in ["nhwc-8h2w32c2w-2d"]: + 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 layout in ["nhwc-4h2w32c2w-2d"]: + N, H, W, C = arr_np.shape + return arr_np.reshape([N, H // 4, 4, W // 4, 2, 2, C // 32, 32]).transpose( + 0, 1, 3, 6, 2, 4, 7, 5 + ) + if layout in ["nc-1024c-2d"]: + N, C = arr_np.shape + return arr_np.reshape([N, C // 1024, 1024]) + if layout in ["nc-512c-2d"]: + N, C = arr_np.shape + return arr_np.reshape([N, C // 512, 512]) + raise RuntimeError(f"Unexpected layout '{layout}'") + + +class TestCastF16F32Slice2d: + input_shape, input_layout, output_layout, axis_sep = tvm.testing.parameters( + ((1, 16, 12, 64), "nhwc-8h2w32c2w-2d", "nhwc-8h2w32c2w-2d", [4]), + ((1, 64, 64, 32), "nhwc-8h2w32c2w-2d", "nhwc-8h2w32c2w-2d", [4]), + ((1, 16, 12, 64), "nhwc-8h2w32c2w-2d", "nhwc-4h2w32c2w-2d", [4]), + ((1, 64, 64, 32), "nhwc-8h2w32c2w-2d", "nhwc-4h2w32c2w-2d", [4]), + ((1, 1024), "nc-1024c-2d", "nc-1024c-2d", [2]), + ((1, 1024), "nc-1024c-2d", "nc-512c-2d", [2]), + ) + dtype = tvm.testing.parameter("float16") + working_scope = tvm.testing.parameter("global.vtcm") + + @tvm.testing.fixture + def input_np(self, input_shape, dtype): + return np.random.uniform(size=input_shape).astype(dtype) + + @tvm.testing.fixture + def transformed_input_np(self, input_np, input_layout): + return transform_numpy(input_np, input_layout) + + @tvm.testing.fixture + def expected_output_np(self, input_np): + ref_np = input_np.astype("float32") + return ref_np + + @tvm.testing.fixture + def transformed_expected_output_np(self, expected_output_np, output_layout): + return transform_numpy(expected_output_np, output_layout) + + @tvm.testing.requires_hexagon + def test_cast_fp16_fp32_slice( + self, + input_shape, + dtype, + input_layout, + output_layout, + input_np, + transformed_input_np, + transformed_expected_output_np, + axis_sep, + hexagon_session, + working_scope, + ): + + target_hexagon = tvm.target.hexagon("v68", llvm_options="--disable-loop-unrolling-pass") + target = tvm.target.Target(target_hexagon, host=target_hexagon) + A = te.placeholder(input_shape, name="A", dtype=dtype) + M = sl.cast_f16_f32_compute(A) + cast_func = te.create_prim_func([A, M]) + tir_s = sl.cast_f16_f32_schedule(cast_func, input_layout, output_layout) + A_data = allocate_hexagon_array( + hexagon_session.device, + data=transformed_input_np, + axis_separators=axis_sep, + mem_scope=working_scope, + ) + M_data = allocate_hexagon_array( + hexagon_session.device, + tensor_shape=transformed_expected_output_np.shape, + dtype=transformed_expected_output_np.dtype, + axis_separators=axis_sep, + mem_scope=working_scope, + ) + with tvm.transform.PassContext(opt_level=3): + tir_irm = tvm.lower(tir_s.mod, [A, M], name="cast_f16_f32") + runtime_module = tvm.build(tir_irm, target=target, name="cast_f16_f32") + mod = hexagon_session.load_module(runtime_module) + + mod(A_data, M_data) + output_np = M_data.numpy() + tvm.testing.assert_allclose( + output_np, + transformed_expected_output_np, + 1e-3, + 1e-3, + ) + + +class TestCastF32F16Slice2d: + (input_shape, input_layout, output_layout, axis_sep,) = tvm.testing.parameters( + ((1, 16, 12, 64), "nhwc-8h2w32c2w-2d", "nhwc-8h2w32c2w-2d", [4]), + ((1, 64, 64, 32), "nhwc-8h2w32c2w-2d", "nhwc-8h2w32c2w-2d", [4]), + ((1, 16, 12, 64), "nhwc-4h2w32c2w-2d", "nhwc-8h2w32c2w-2d", [4]), + ((1, 64, 64, 32), "nhwc-4h2w32c2w-2d", "nhwc-8h2w32c2w-2d", [4]), + ((1, 1024), "nc-1024c-2d", "nc-1024c-2d", [2]), + ((1, 1024), "nc-512c-2d", "nc-1024c-2d", [2]), + ) + dtype = tvm.testing.parameter("float32") + working_scope = tvm.testing.parameter("global.vtcm") + + @tvm.testing.fixture + def input_np(self, input_shape, dtype): + return np.random.uniform(size=input_shape).astype(dtype) + + @tvm.testing.fixture + def transformed_input_np(self, input_np, input_layout): + return transform_numpy(input_np, input_layout) + + @tvm.testing.fixture + def expected_output_np(self, input_np): + ref_np = input_np.astype("float16") + return ref_np + + @tvm.testing.fixture + def transformed_expected_output_np(self, expected_output_np, output_layout): + return transform_numpy(expected_output_np, output_layout) + + @tvm.testing.requires_hexagon + def test_cast_fp32_fp16_slice( + self, + input_shape, + dtype, + input_layout, + output_layout, + input_np, + transformed_input_np, + transformed_expected_output_np, + axis_sep, + hexagon_session, + working_scope, + ): + + target_hexagon = tvm.target.hexagon("v68", llvm_options="--disable-loop-unrolling-pass") + target = tvm.target.Target(target_hexagon, host=target_hexagon) + A = te.placeholder(input_shape, name="A", dtype=dtype) + M = sl.cast_f32_f16_compute(A) + cast_func = te.create_prim_func([A, M]) + tir_s = sl.cast_f32_f16_schedule(cast_func, input_layout, output_layout) + A_data = allocate_hexagon_array( + hexagon_session.device, + data=transformed_input_np, + axis_separators=axis_sep, + mem_scope=working_scope, + ) + M_data = allocate_hexagon_array( + hexagon_session.device, + tensor_shape=transformed_expected_output_np.shape, + dtype=transformed_expected_output_np.dtype, + axis_separators=axis_sep, + mem_scope=working_scope, + ) + with tvm.transform.PassContext(opt_level=3): + tir_irm = tvm.lower(tir_s.mod, [A, M], name="cast_f32_f16") + runtime_module = tvm.build(tir_irm, target=target, name="cast_f32_f16") + mod = hexagon_session.load_module(runtime_module) + + mod(A_data, M_data) + output_np = M_data.numpy() + tvm.testing.assert_allclose( + output_np, + transformed_expected_output_np, + 1e-3, + 1e-3, + ) + + +if __name__ == "__main__": + sys.exit(pytest.main(sys.argv)) From 7b8b1d20e399490fb74840b80784f7a2ca952114 Mon Sep 17 00:00:00 2001 From: Arun Rangasamy Date: Wed, 15 Jun 2022 17:06:09 +0530 Subject: [PATCH 02/11] Fix linter issues. Move test to topi directory. --- .../{ => topi}/test_cast_slice.py | 28 +++++++++++++------ 1 file changed, 19 insertions(+), 9 deletions(-) rename tests/python/contrib/test_hexagon/{ => topi}/test_cast_slice.py (93%) diff --git a/tests/python/contrib/test_hexagon/test_cast_slice.py b/tests/python/contrib/test_hexagon/topi/test_cast_slice.py similarity index 93% rename from tests/python/contrib/test_hexagon/test_cast_slice.py rename to tests/python/contrib/test_hexagon/topi/test_cast_slice.py index 67cb3ecbde79..9c012a3bc544 100644 --- a/tests/python/contrib/test_hexagon/test_cast_slice.py +++ b/tests/python/contrib/test_hexagon/topi/test_cast_slice.py @@ -20,16 +20,17 @@ import tvm import tvm.testing -from tvm import te, topi +from tvm import te import tvm.topi.hexagon.slice_ops as sl -from tvm.topi import testing -from .infrastructure import allocate_hexagon_array -import tvm.contrib.hexagon +from ..infrastructure import allocate_hexagon_array # pylint: disable=invalid-name def transform_numpy(arr_np, layout): + """ + Layout transformation on numpy arrays + """ if layout in ["nhwc-8h2w32c2w-2d"]: N, H, W, C = arr_np.shape return arr_np.reshape([N, H // 8, 8, W // 4, 2, 2, C // 32, 32]).transpose( @@ -50,6 +51,9 @@ def transform_numpy(arr_np, layout): class TestCastF16F32Slice2d: + """ + For testing Cast F16 to F32 Slice ops + """ input_shape, input_layout, output_layout, axis_sep = tvm.testing.parameters( ((1, 16, 12, 64), "nhwc-8h2w32c2w-2d", "nhwc-8h2w32c2w-2d", [4]), ((1, 64, 64, 32), "nhwc-8h2w32c2w-2d", "nhwc-8h2w32c2w-2d", [4]), @@ -85,15 +89,16 @@ def test_cast_fp16_fp32_slice( dtype, input_layout, output_layout, - input_np, transformed_input_np, transformed_expected_output_np, axis_sep, hexagon_session, working_scope, ): - - target_hexagon = tvm.target.hexagon("v68", llvm_options="--disable-loop-unrolling-pass") + """ + Top level testing function for cast fp16 to fp32 + """ + target_hexagon = tvm.target.hexagon("v68") target = tvm.target.Target(target_hexagon, host=target_hexagon) A = te.placeholder(input_shape, name="A", dtype=dtype) M = sl.cast_f16_f32_compute(A) @@ -128,6 +133,9 @@ def test_cast_fp16_fp32_slice( class TestCastF32F16Slice2d: + """ + For testing Cast F32 to F16 Slice ops + """ (input_shape, input_layout, output_layout, axis_sep,) = tvm.testing.parameters( ((1, 16, 12, 64), "nhwc-8h2w32c2w-2d", "nhwc-8h2w32c2w-2d", [4]), ((1, 64, 64, 32), "nhwc-8h2w32c2w-2d", "nhwc-8h2w32c2w-2d", [4]), @@ -163,15 +171,17 @@ def test_cast_fp32_fp16_slice( dtype, input_layout, output_layout, - input_np, transformed_input_np, transformed_expected_output_np, axis_sep, hexagon_session, working_scope, ): + """ + Top level testing function for cast fp32 to fp16 + """ - target_hexagon = tvm.target.hexagon("v68", llvm_options="--disable-loop-unrolling-pass") + target_hexagon = tvm.target.hexagon("v68") target = tvm.target.Target(target_hexagon, host=target_hexagon) A = te.placeholder(input_shape, name="A", dtype=dtype) M = sl.cast_f32_f16_compute(A) From 7d710e5a7f7b647584350d14240b4220c965050a Mon Sep 17 00:00:00 2001 From: Arun Rangasamy Date: Wed, 15 Jun 2022 17:10:08 +0530 Subject: [PATCH 03/11] run through black. --- tests/python/contrib/test_hexagon/topi/test_cast_slice.py | 2 ++ 1 file changed, 2 insertions(+) diff --git a/tests/python/contrib/test_hexagon/topi/test_cast_slice.py b/tests/python/contrib/test_hexagon/topi/test_cast_slice.py index 9c012a3bc544..4bef0b4f8d10 100644 --- a/tests/python/contrib/test_hexagon/topi/test_cast_slice.py +++ b/tests/python/contrib/test_hexagon/topi/test_cast_slice.py @@ -54,6 +54,7 @@ class TestCastF16F32Slice2d: """ For testing Cast F16 to F32 Slice ops """ + input_shape, input_layout, output_layout, axis_sep = tvm.testing.parameters( ((1, 16, 12, 64), "nhwc-8h2w32c2w-2d", "nhwc-8h2w32c2w-2d", [4]), ((1, 64, 64, 32), "nhwc-8h2w32c2w-2d", "nhwc-8h2w32c2w-2d", [4]), @@ -136,6 +137,7 @@ class TestCastF32F16Slice2d: """ For testing Cast F32 to F16 Slice ops """ + (input_shape, input_layout, output_layout, axis_sep,) = tvm.testing.parameters( ((1, 16, 12, 64), "nhwc-8h2w32c2w-2d", "nhwc-8h2w32c2w-2d", [4]), ((1, 64, 64, 32), "nhwc-8h2w32c2w-2d", "nhwc-8h2w32c2w-2d", [4]), From c7eb1622ca1ab75d86820fa7fb9b407bc3285e70 Mon Sep 17 00:00:00 2001 From: Arun Rangasamy Date: Thu, 16 Jun 2022 15:46:17 +0530 Subject: [PATCH 04/11] Fix pylint error --- python/tvm/topi/hexagon/slice_ops/__init__.py | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/python/tvm/topi/hexagon/slice_ops/__init__.py b/python/tvm/topi/hexagon/slice_ops/__init__.py index 36f10e2e0aab..93b7f4377580 100644 --- a/python/tvm/topi/hexagon/slice_ops/__init__.py +++ b/python/tvm/topi/hexagon/slice_ops/__init__.py @@ -20,4 +20,9 @@ # pylint: disable=wildcard-import from .avg_pool2d import avg_pool2d_compute, avg_pool2d_STIR_schedule -from .cast import cast_f16_f32_compute, cast_f16_f32_schedule, cast_f32_f16_compute, cast_f32_f16_schedule +from .cast import ( + cast_f16_f32_compute, + cast_f16_f32_schedule, + cast_f32_f16_compute, + cast_f32_f16_schedule, +) From 5e98f1bd18a19138068f93e10cccf39d603515a8 Mon Sep 17 00:00:00 2001 From: Arun Rangasamy Date: Wed, 22 Jun 2022 10:37:13 +0530 Subject: [PATCH 05/11] Use tvm.testing.main --- tests/python/contrib/test_hexagon/topi/test_cast_slice.py | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/tests/python/contrib/test_hexagon/topi/test_cast_slice.py b/tests/python/contrib/test_hexagon/topi/test_cast_slice.py index 4bef0b4f8d10..394db90d3107 100644 --- a/tests/python/contrib/test_hexagon/topi/test_cast_slice.py +++ b/tests/python/contrib/test_hexagon/topi/test_cast_slice.py @@ -99,7 +99,7 @@ def test_cast_fp16_fp32_slice( """ Top level testing function for cast fp16 to fp32 """ - target_hexagon = tvm.target.hexagon("v68") + target_hexagon = tvm.target.hexagon("v69") target = tvm.target.Target(target_hexagon, host=target_hexagon) A = te.placeholder(input_shape, name="A", dtype=dtype) M = sl.cast_f16_f32_compute(A) @@ -183,7 +183,7 @@ def test_cast_fp32_fp16_slice( Top level testing function for cast fp32 to fp16 """ - target_hexagon = tvm.target.hexagon("v68") + target_hexagon = tvm.target.hexagon("v69") target = tvm.target.Target(target_hexagon, host=target_hexagon) A = te.placeholder(input_shape, name="A", dtype=dtype) M = sl.cast_f32_f16_compute(A) @@ -218,4 +218,4 @@ def test_cast_fp32_fp16_slice( if __name__ == "__main__": - sys.exit(pytest.main(sys.argv)) + sys.exit(tvm.testing.main(sys.argv)) From f531e0046a3d9cca5dfe3c8c2034c24144753425 Mon Sep 17 00:00:00 2001 From: Arun Rangasamy Date: Thu, 23 Jun 2022 12:48:03 +0530 Subject: [PATCH 06/11] Reuse functions --- python/tvm/topi/hexagon/slice_ops/cast.py | 69 ++++------------- python/tvm/topi/hexagon/utils.py | 21 ++++++ .../contrib/test_hexagon/infrastructure.py | 15 +++- .../test_hexagon/topi/test_cast_slice.py | 74 +++++++------------ 4 files changed, 72 insertions(+), 107 deletions(-) diff --git a/python/tvm/topi/hexagon/slice_ops/cast.py b/python/tvm/topi/hexagon/slice_ops/cast.py index 3e8024da733a..4889b11444cc 100644 --- a/python/tvm/topi/hexagon/slice_ops/cast.py +++ b/python/tvm/topi/hexagon/slice_ops/cast.py @@ -18,70 +18,25 @@ from tvm import te from tvm import tir -from tvm.tir import IndexMap +from ..utils import get_layout_transform_fn # pylint: disable=invalid-name -def layout_transform_nhwc_8h2w32c2w(n, h, w, c): - return [ - n, - h // 8, - w // 4, - c // 32, - IndexMap.AXIS_SEPARATOR, - h % 8, - (w % 4) // 2, - c % 32, - w % 2, - ] - - -def layout_transform_nc_1024c(n, c): - return [ - n, - c // 1024, - IndexMap.AXIS_SEPARATOR, - c % 1024, - ] - - -def layout_transform_nhwc_4h2w32c2w(n, h, w, c): - return [ - n, - h // 4, - w // 4, - c // 32, - IndexMap.AXIS_SEPARATOR, - h % 4, - (w % 4) // 2, - (c % 32), - w % 2, - ] - - -def layout_transform_nc_512c(n, c): - return [ - n, - c // 512, - IndexMap.AXIS_SEPARATOR, - c % 512, - ] - - def get_layout_transform_for_f32(f32_layout_string): """ Given f32 layout string, return transform_layout function and channel/height split factor to be used for scheduling """ + layout_transform_fn = get_layout_transform_fn(f32_layout_string) if f32_layout_string == "nhwc-8h2w32c2w-2d": - return [layout_transform_nhwc_8h2w32c2w, 8] + return [layout_transform_fn, 8] if f32_layout_string == "nhwc-4h2w32c2w-2d": - return [layout_transform_nhwc_4h2w32c2w, 4] + return [layout_transform_fn, 4] if f32_layout_string == "nc-1024c-2d": - return [layout_transform_nc_1024c, 1024] + return [layout_transform_fn, 1024] if f32_layout_string == "nc-512c-2d": - return [layout_transform_nc_512c, 512] + return [layout_transform_fn, 512] raise RuntimeError(f"Unexpected f32_layout '{f32_layout_string}'") @@ -124,16 +79,17 @@ def cast_f16_f32_stir_schedule_nc(func, in_layout, out_layout, c_split_factor): def cast_f16_f32_schedule(cast_func, in_layout_str, out_layout_str): """Schedule for f16 to f32 cast: top level function""" f32_layout_transform_func, split_factor = get_layout_transform_for_f32(out_layout_str) + f16_layout_transform_func = get_layout_transform_fn(in_layout_str) if in_layout_str == "nhwc-8h2w32c2w-2d": return cast_f16_f32_stir_schedule_nhwc( cast_func, - layout_transform_nhwc_8h2w32c2w, + f16_layout_transform_func, f32_layout_transform_func, split_factor, ) if in_layout_str == "nc-1024c-2d": return cast_f16_f32_stir_schedule_nc( - cast_func, layout_transform_nc_1024c, f32_layout_transform_func, split_factor + cast_func, f16_layout_transform_func, f32_layout_transform_func, split_factor ) raise RuntimeError(f"Unexpected input_layout, output_layout '{input_layout, output_layout}'") @@ -177,12 +133,13 @@ def cast_f32_f16_stir_schedule_nc(func, in_layout, out_layout, c_split_factor): def cast_f32_f16_schedule(cast_func, in_layout_str, out_layout_str): """Schedule for f32 to f16 cast: top level function""" f32_layout_transform_func, split_factor = get_layout_transform_for_f32(in_layout_str) + f16_layout_transform_func = get_layout_transform_fn(out_layout_str) if out_layout_str == "nhwc-8h2w32c2w-2d": return cast_f32_f16_stir_schedule_nhwc( - cast_func, f32_layout_transform_func, layout_transform_nhwc_8h2w32c2w, split_factor + cast_func, f32_layout_transform_func, f16_layout_transform_func, split_factor ) if out_layout_str == "nc-1024c-2d": return cast_f32_f16_stir_schedule_nc( - cast_func, f32_layout_transform_func, layout_transform_nc_1024c, split_factor + cast_func, f32_layout_transform_func, f16_layout_transform_func, split_factor ) - raise RuntimeError(f"Unexpected input_layout, output_layout '{input_layout, output_layout}'") + raise RuntimeError(f"Unexpected input_layout, output_layout '{in_layout_str, out_layout_str}'") diff --git a/python/tvm/topi/hexagon/utils.py b/python/tvm/topi/hexagon/utils.py index af6e3de9c350..68b8d64ef0ae 100644 --- a/python/tvm/topi/hexagon/utils.py +++ b/python/tvm/topi/hexagon/utils.py @@ -39,6 +39,21 @@ 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 nc_1024c(n, c): + """Return index map for nc_1024c 2d layout""" + return [n, c // 1024, te.AXIS_SEPARATOR, c % 1024] + + +def nhwc_4h2w32c2w_2d(n, h, w, c): + """Return index map for nhwc_4h2w32c2w 2d layout""" + return [n, h // 4, w // 4, c // 32, te.AXIS_SEPARATOR, h % 4, (w % 4) // 2, c % 32, w % 2] + + +def nc_512c(n, c): + """Return index map for nc_512c 2d layout""" + return [n, c // 512, te.AXIS_SEPARATOR, c % 512] + + def get_layout_transform_fn(layout): """Return index map function as per the layout string""" if layout == "nhwc-8h2w32c2w-2d": @@ -49,4 +64,10 @@ def get_layout_transform_fn(layout): return n11c_1024c_2d if layout == "n11c-1024c-1d": return n11c_1024c_1d + if layout == "nhwc-4h2w32c2w-2d": + return nhwc_4h2w32c2w_2d + if layout == "nc-1024c-2d": + return nc_1024c + if layout == "nc-512c-2d": + return nc_512c 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 57a9dff8b424..07977d474802 100644 --- a/tests/python/contrib/test_hexagon/infrastructure.py +++ b/tests/python/contrib/test_hexagon/infrastructure.py @@ -234,17 +234,28 @@ def compute(n, ho, wo, ko, hi, wi, ki): 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": + n, h, w, c = arr_np.shape 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 ["nhwc-4h2w32c2w-2d"]: + return arr_np.reshape([n, h // 4, 4, 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}'") + if current_layout == "nc": + n, c = arr_np.shape + if new_layout in ["nc-1024c-2d"]: + return arr_np.reshape([n, c // 1024, 1024]) + if new_layout in ["nc-512c-2d"]: + return arr_np.reshape([n, c // 512, 512]) raise RuntimeError(f"Unexpected new_layout '{new_layout}'") + raise RuntimeError(f"Unexpected current_layout '{current_layout}'") diff --git a/tests/python/contrib/test_hexagon/topi/test_cast_slice.py b/tests/python/contrib/test_hexagon/topi/test_cast_slice.py index 394db90d3107..725133c0dca5 100644 --- a/tests/python/contrib/test_hexagon/topi/test_cast_slice.py +++ b/tests/python/contrib/test_hexagon/topi/test_cast_slice.py @@ -16,52 +16,28 @@ # under the License. """ Tests for Hexagon slice cast ops """ import numpy as np -import pytest import tvm import tvm.testing from tvm import te import tvm.topi.hexagon.slice_ops as sl -from ..infrastructure import allocate_hexagon_array +from ..infrastructure import allocate_hexagon_array, transform_numpy # pylint: disable=invalid-name -def transform_numpy(arr_np, layout): - """ - Layout transformation on numpy arrays - """ - if layout in ["nhwc-8h2w32c2w-2d"]: - 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 layout in ["nhwc-4h2w32c2w-2d"]: - N, H, W, C = arr_np.shape - return arr_np.reshape([N, H // 4, 4, W // 4, 2, 2, C // 32, 32]).transpose( - 0, 1, 3, 6, 2, 4, 7, 5 - ) - if layout in ["nc-1024c-2d"]: - N, C = arr_np.shape - return arr_np.reshape([N, C // 1024, 1024]) - if layout in ["nc-512c-2d"]: - N, C = arr_np.shape - return arr_np.reshape([N, C // 512, 512]) - raise RuntimeError(f"Unexpected layout '{layout}'") - - class TestCastF16F32Slice2d: """ For testing Cast F16 to F32 Slice ops """ - input_shape, input_layout, output_layout, axis_sep = tvm.testing.parameters( - ((1, 16, 12, 64), "nhwc-8h2w32c2w-2d", "nhwc-8h2w32c2w-2d", [4]), - ((1, 64, 64, 32), "nhwc-8h2w32c2w-2d", "nhwc-8h2w32c2w-2d", [4]), - ((1, 16, 12, 64), "nhwc-8h2w32c2w-2d", "nhwc-4h2w32c2w-2d", [4]), - ((1, 64, 64, 32), "nhwc-8h2w32c2w-2d", "nhwc-4h2w32c2w-2d", [4]), - ((1, 1024), "nc-1024c-2d", "nc-1024c-2d", [2]), - ((1, 1024), "nc-1024c-2d", "nc-512c-2d", [2]), + input_shape, orig_layout, input_layout, output_layout, axis_sep = tvm.testing.parameters( + ((1, 16, 12, 64), "nhwc", "nhwc-8h2w32c2w-2d", "nhwc-8h2w32c2w-2d", [4]), + ((1, 64, 64, 32), "nhwc", "nhwc-8h2w32c2w-2d", "nhwc-8h2w32c2w-2d", [4]), + ((1, 16, 12, 64), "nhwc", "nhwc-8h2w32c2w-2d", "nhwc-4h2w32c2w-2d", [4]), + ((1, 64, 64, 32), "nhwc", "nhwc-8h2w32c2w-2d", "nhwc-4h2w32c2w-2d", [4]), + ((1, 1024), "nc", "nc-1024c-2d", "nc-1024c-2d", [2]), + ((1, 1024), "nc", "nc-1024c-2d", "nc-512c-2d", [2]), ) dtype = tvm.testing.parameter("float16") working_scope = tvm.testing.parameter("global.vtcm") @@ -71,8 +47,8 @@ def input_np(self, input_shape, dtype): return np.random.uniform(size=input_shape).astype(dtype) @tvm.testing.fixture - def transformed_input_np(self, input_np, input_layout): - return transform_numpy(input_np, input_layout) + def transformed_input_np(self, input_np, orig_layout, input_layout): + return transform_numpy(input_np, orig_layout, input_layout) @tvm.testing.fixture def expected_output_np(self, input_np): @@ -80,8 +56,8 @@ def expected_output_np(self, input_np): return ref_np @tvm.testing.fixture - def transformed_expected_output_np(self, expected_output_np, output_layout): - return transform_numpy(expected_output_np, output_layout) + def transformed_expected_output_np(self, expected_output_np, orig_layout, output_layout): + return transform_numpy(expected_output_np, orig_layout, output_layout) @tvm.testing.requires_hexagon def test_cast_fp16_fp32_slice( @@ -99,7 +75,7 @@ def test_cast_fp16_fp32_slice( """ Top level testing function for cast fp16 to fp32 """ - target_hexagon = tvm.target.hexagon("v69") + target_hexagon = tvm.target.hexagon("v68") target = tvm.target.Target(target_hexagon, host=target_hexagon) A = te.placeholder(input_shape, name="A", dtype=dtype) M = sl.cast_f16_f32_compute(A) @@ -138,13 +114,13 @@ class TestCastF32F16Slice2d: For testing Cast F32 to F16 Slice ops """ - (input_shape, input_layout, output_layout, axis_sep,) = tvm.testing.parameters( - ((1, 16, 12, 64), "nhwc-8h2w32c2w-2d", "nhwc-8h2w32c2w-2d", [4]), - ((1, 64, 64, 32), "nhwc-8h2w32c2w-2d", "nhwc-8h2w32c2w-2d", [4]), - ((1, 16, 12, 64), "nhwc-4h2w32c2w-2d", "nhwc-8h2w32c2w-2d", [4]), - ((1, 64, 64, 32), "nhwc-4h2w32c2w-2d", "nhwc-8h2w32c2w-2d", [4]), - ((1, 1024), "nc-1024c-2d", "nc-1024c-2d", [2]), - ((1, 1024), "nc-512c-2d", "nc-1024c-2d", [2]), + (input_shape, orig_layout, input_layout, output_layout, axis_sep,) = tvm.testing.parameters( + ((1, 16, 12, 64), "nhwc", "nhwc-8h2w32c2w-2d", "nhwc-8h2w32c2w-2d", [4]), + ((1, 64, 64, 32), "nhwc", "nhwc-8h2w32c2w-2d", "nhwc-8h2w32c2w-2d", [4]), + ((1, 16, 12, 64), "nhwc", "nhwc-4h2w32c2w-2d", "nhwc-8h2w32c2w-2d", [4]), + ((1, 64, 64, 32), "nhwc", "nhwc-4h2w32c2w-2d", "nhwc-8h2w32c2w-2d", [4]), + ((1, 1024), "nc", "nc-1024c-2d", "nc-1024c-2d", [2]), + ((1, 1024), "nc", "nc-512c-2d", "nc-1024c-2d", [2]), ) dtype = tvm.testing.parameter("float32") working_scope = tvm.testing.parameter("global.vtcm") @@ -154,8 +130,8 @@ def input_np(self, input_shape, dtype): return np.random.uniform(size=input_shape).astype(dtype) @tvm.testing.fixture - def transformed_input_np(self, input_np, input_layout): - return transform_numpy(input_np, input_layout) + def transformed_input_np(self, input_np, orig_layout, input_layout): + return transform_numpy(input_np, orig_layout, input_layout) @tvm.testing.fixture def expected_output_np(self, input_np): @@ -163,8 +139,8 @@ def expected_output_np(self, input_np): return ref_np @tvm.testing.fixture - def transformed_expected_output_np(self, expected_output_np, output_layout): - return transform_numpy(expected_output_np, output_layout) + def transformed_expected_output_np(self, expected_output_np, orig_layout, output_layout): + return transform_numpy(expected_output_np, orig_layout, output_layout) @tvm.testing.requires_hexagon def test_cast_fp32_fp16_slice( @@ -183,7 +159,7 @@ def test_cast_fp32_fp16_slice( Top level testing function for cast fp32 to fp16 """ - target_hexagon = tvm.target.hexagon("v69") + target_hexagon = tvm.target.hexagon("v68") target = tvm.target.Target(target_hexagon, host=target_hexagon) A = te.placeholder(input_shape, name="A", dtype=dtype) M = sl.cast_f32_f16_compute(A) From 258828c2ff609a7f4b65fa7e3839c94a36134aa1 Mon Sep 17 00:00:00 2001 From: Arun Rangasamy Date: Sun, 26 Jun 2022 08:23:36 +0530 Subject: [PATCH 07/11] Fix tvm.testing.main invocation --- tests/python/contrib/test_hexagon/topi/test_cast_slice.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/python/contrib/test_hexagon/topi/test_cast_slice.py b/tests/python/contrib/test_hexagon/topi/test_cast_slice.py index 725133c0dca5..49f1835b7b73 100644 --- a/tests/python/contrib/test_hexagon/topi/test_cast_slice.py +++ b/tests/python/contrib/test_hexagon/topi/test_cast_slice.py @@ -194,4 +194,4 @@ def test_cast_fp32_fp16_slice( if __name__ == "__main__": - sys.exit(tvm.testing.main(sys.argv)) + sys.exit(tvm.testing.main()) From 2a3fb2a66110261c3aedd64fbfd209330422b128 Mon Sep 17 00:00:00 2001 From: Arun Rangasamy Date: Tue, 28 Jun 2022 11:00:13 +0530 Subject: [PATCH 08/11] Fix lint issues; Address review comments --- python/tvm/topi/hexagon/slice_ops/cast.py | 30 ++++++++-------- .../test_hexagon/topi/test_cast_slice.py | 36 +++++++++---------- 2 files changed, 31 insertions(+), 35 deletions(-) diff --git a/python/tvm/topi/hexagon/slice_ops/cast.py b/python/tvm/topi/hexagon/slice_ops/cast.py index 4889b11444cc..b4984763e0e0 100644 --- a/python/tvm/topi/hexagon/slice_ops/cast.py +++ b/python/tvm/topi/hexagon/slice_ops/cast.py @@ -20,8 +20,6 @@ from tvm import tir from ..utils import get_layout_transform_fn -# pylint: disable=invalid-name - def get_layout_transform_for_f32(f32_layout_string): """ @@ -51,12 +49,12 @@ def cast_f16_f32_stir_schedule_nhwc(func, in_layout, out_layout, h_split_factor) """Schedule for nhwc f16 to f32 cast: nhwc layout""" sch = tir.Schedule(func, debug_mask="all") block_name = "CastF16F32" - n, h, w, c = sch.get_loops(sch.get_block(block_name)) - h_outer, h_inner = sch.split(h, [None, h_split_factor]) - w_outer, w_inner = sch.split(w, [None, 4]) - c_outer, c_inner = sch.split(c, [None, 32]) + n_orig, h_orig, w_orig, c_orig = sch.get_loops(sch.get_block(block_name)) + h_outer, h_inner = sch.split(h_orig, [None, h_split_factor]) + w_outer, w_inner = sch.split(w_orig, [None, 4]) + c_outer, c_inner = sch.split(c_orig, [None, 32]) w_inner_o, w_inner_i = sch.split(w_inner, [None, 2]) - sch.reorder(n, h_outer, w_outer, c_outer, h_inner, w_inner_o, c_inner, w_inner_i) + sch.reorder(n_orig, h_outer, w_outer, c_outer, h_inner, w_inner_o, c_inner, w_inner_i) sch.transform_layout(block_name, "A", in_layout) sch.transform_layout(block_name, block_name, out_layout) fused = sch.fuse(c_inner, w_inner_i) @@ -68,8 +66,8 @@ def cast_f16_f32_stir_schedule_nc(func, in_layout, out_layout, c_split_factor): """Schedule for nc f16 to f32 cast: nc layout""" sch = tir.Schedule(func, debug_mask="all") block_name = "CastF16F32" - _, c = sch.get_loops(sch.get_block(block_name)) - _, c_inner = sch.split(c, [None, c_split_factor]) + _, c_orig = sch.get_loops(sch.get_block(block_name)) + _, c_inner = sch.split(c_orig, [None, c_split_factor]) sch.transform_layout(block_name, "A", in_layout) sch.transform_layout(block_name, block_name, out_layout) sch.vectorize(c_inner) @@ -105,12 +103,12 @@ def cast_f32_f16_stir_schedule_nhwc(func, in_layout, out_layout, h_split_factor) """Schedule for nhwc f32 to f16 cast: nhwc layout""" sch = tir.Schedule(func, debug_mask="all") block_name = "CastF32F16" - n, h, w, c = sch.get_loops(sch.get_block(block_name)) - h_outer, h_inner = sch.split(h, [None, h_split_factor]) - w_outer, w_inner = sch.split(w, [None, 4]) - c_outer, c_inner = sch.split(c, [None, 32]) + n_orig, h_orig, w_orig, c_orig = sch.get_loops(sch.get_block(block_name)) + h_outer, h_inner = sch.split(h_orig, [None, h_split_factor]) + w_outer, w_inner = sch.split(w_orig, [None, 4]) + c_outer, c_inner = sch.split(c_orig, [None, 32]) w_inner_o, w_inner_i = sch.split(w_inner, [None, 2]) - sch.reorder(n, h_outer, w_outer, c_outer, h_inner, w_inner_o, c_inner, w_inner_i) + sch.reorder(n_orig, h_outer, w_outer, c_outer, h_inner, w_inner_o, c_inner, w_inner_i) sch.transform_layout(block_name, "A", in_layout) sch.transform_layout(block_name, block_name, out_layout) fused = sch.fuse(c_inner, w_inner_i) @@ -122,8 +120,8 @@ def cast_f32_f16_stir_schedule_nc(func, in_layout, out_layout, c_split_factor): """Schedule for nc f32 to f16 cast: nc layout""" sch = tir.Schedule(func, debug_mask="all") block_name = "CastF32F16" - _, c = sch.get_loops(sch.get_block(block_name)) - _, c_inner = sch.split(c, [None, c_split_factor]) + _, c_orig = sch.get_loops(sch.get_block(block_name)) + _, c_inner = sch.split(c_orig, [None, c_split_factor]) sch.transform_layout(block_name, "A", in_layout) sch.transform_layout(block_name, block_name, out_layout) sch.vectorize(c_inner) diff --git a/tests/python/contrib/test_hexagon/topi/test_cast_slice.py b/tests/python/contrib/test_hexagon/topi/test_cast_slice.py index 49f1835b7b73..3407091ad64b 100644 --- a/tests/python/contrib/test_hexagon/topi/test_cast_slice.py +++ b/tests/python/contrib/test_hexagon/topi/test_cast_slice.py @@ -23,8 +23,6 @@ import tvm.topi.hexagon.slice_ops as sl from ..infrastructure import allocate_hexagon_array, transform_numpy -# pylint: disable=invalid-name - class TestCastF16F32Slice2d: """ @@ -77,17 +75,17 @@ def test_cast_fp16_fp32_slice( """ target_hexagon = tvm.target.hexagon("v68") target = tvm.target.Target(target_hexagon, host=target_hexagon) - A = te.placeholder(input_shape, name="A", dtype=dtype) - M = sl.cast_f16_f32_compute(A) - cast_func = te.create_prim_func([A, M]) + cast_input = te.placeholder(input_shape, name="A", dtype=dtype) + cast_output = sl.cast_f16_f32_compute(cast_input) + cast_func = te.create_prim_func([cast_input, cast_output]) tir_s = sl.cast_f16_f32_schedule(cast_func, input_layout, output_layout) - A_data = allocate_hexagon_array( + input_data = allocate_hexagon_array( hexagon_session.device, data=transformed_input_np, axis_separators=axis_sep, mem_scope=working_scope, ) - M_data = allocate_hexagon_array( + output_data = allocate_hexagon_array( hexagon_session.device, tensor_shape=transformed_expected_output_np.shape, dtype=transformed_expected_output_np.dtype, @@ -95,12 +93,12 @@ def test_cast_fp16_fp32_slice( mem_scope=working_scope, ) with tvm.transform.PassContext(opt_level=3): - tir_irm = tvm.lower(tir_s.mod, [A, M], name="cast_f16_f32") + tir_irm = tvm.lower(tir_s.mod, [cast_input, cast_output], name="cast_f16_f32") runtime_module = tvm.build(tir_irm, target=target, name="cast_f16_f32") mod = hexagon_session.load_module(runtime_module) - mod(A_data, M_data) - output_np = M_data.numpy() + mod(input_data, output_data) + output_np = output_data.numpy() tvm.testing.assert_allclose( output_np, transformed_expected_output_np, @@ -161,17 +159,17 @@ def test_cast_fp32_fp16_slice( target_hexagon = tvm.target.hexagon("v68") target = tvm.target.Target(target_hexagon, host=target_hexagon) - A = te.placeholder(input_shape, name="A", dtype=dtype) - M = sl.cast_f32_f16_compute(A) - cast_func = te.create_prim_func([A, M]) + cast_input = te.placeholder(input_shape, name="A", dtype=dtype) + cast_output = sl.cast_f32_f16_compute(cast_input) + cast_func = te.create_prim_func([cast_input, cast_output]) tir_s = sl.cast_f32_f16_schedule(cast_func, input_layout, output_layout) - A_data = allocate_hexagon_array( + input_data = allocate_hexagon_array( hexagon_session.device, data=transformed_input_np, axis_separators=axis_sep, mem_scope=working_scope, ) - M_data = allocate_hexagon_array( + output_data = allocate_hexagon_array( hexagon_session.device, tensor_shape=transformed_expected_output_np.shape, dtype=transformed_expected_output_np.dtype, @@ -179,12 +177,12 @@ def test_cast_fp32_fp16_slice( mem_scope=working_scope, ) with tvm.transform.PassContext(opt_level=3): - tir_irm = tvm.lower(tir_s.mod, [A, M], name="cast_f32_f16") + tir_irm = tvm.lower(tir_s.mod, [cast_input, cast_output], name="cast_f32_f16") runtime_module = tvm.build(tir_irm, target=target, name="cast_f32_f16") mod = hexagon_session.load_module(runtime_module) - mod(A_data, M_data) - output_np = M_data.numpy() + mod(input_data, output_data) + output_np = output_data.numpy() tvm.testing.assert_allclose( output_np, transformed_expected_output_np, @@ -194,4 +192,4 @@ def test_cast_fp32_fp16_slice( if __name__ == "__main__": - sys.exit(tvm.testing.main()) + tvm.testing.main() From 7a82e87faea2ed149aff6285387ec6315372217f Mon Sep 17 00:00:00 2001 From: Arun Rangasamy Date: Thu, 30 Jun 2022 10:42:19 +0530 Subject: [PATCH 09/11] run through black. --- python/tvm/topi/hexagon/utils.py | 1 + 1 file changed, 1 insertion(+) diff --git a/python/tvm/topi/hexagon/utils.py b/python/tvm/topi/hexagon/utils.py index 46cd006f0f1e..96c8323ed660 100644 --- a/python/tvm/topi/hexagon/utils.py +++ b/python/tvm/topi/hexagon/utils.py @@ -66,6 +66,7 @@ def nc_512c_2d(n, c): """Return index map for nc_512c 2d layout""" return [n, c // 512, te.AXIS_SEPARATOR, c % 512] + def nc_1024c_2d(n, c): """Return index map for nc_1024c 2d layout""" return [n, c // 1024, te.AXIS_SEPARATOR, c % 1024] From 1196853b56bc73daff50d3107c64aa4e92afb60d Mon Sep 17 00:00:00 2001 From: Arun Rangasamy Date: Mon, 4 Jul 2022 17:46:23 +0530 Subject: [PATCH 10/11] Disable tests on hardware. --- tests/python/contrib/test_hexagon/topi/test_cast_slice.py | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/tests/python/contrib/test_hexagon/topi/test_cast_slice.py b/tests/python/contrib/test_hexagon/topi/test_cast_slice.py index 3407091ad64b..30ea4c94b8b1 100644 --- a/tests/python/contrib/test_hexagon/topi/test_cast_slice.py +++ b/tests/python/contrib/test_hexagon/topi/test_cast_slice.py @@ -73,6 +73,8 @@ def test_cast_fp16_fp32_slice( """ Top level testing function for cast fp16 to fp32 """ + if hexagon_session._launcher._serial_number != "simulator": + pytest.skip(msg="Due to https://github.com/apache/tvm/issues/11957") target_hexagon = tvm.target.hexagon("v68") target = tvm.target.Target(target_hexagon, host=target_hexagon) cast_input = te.placeholder(input_shape, name="A", dtype=dtype) @@ -156,6 +158,8 @@ def test_cast_fp32_fp16_slice( """ Top level testing function for cast fp32 to fp16 """ + if hexagon_session._launcher._serial_number != "simulator": + pytest.skip(msg="Due to https://github.com/apache/tvm/issues/11957") target_hexagon = tvm.target.hexagon("v68") target = tvm.target.Target(target_hexagon, host=target_hexagon) From 6c8392ab4a3c4d6362304c444469e87f6cf5ed8a Mon Sep 17 00:00:00 2001 From: Arun Rangasamy Date: Mon, 4 Jul 2022 21:27:28 +0530 Subject: [PATCH 11/11] Fix test failure --- tests/python/contrib/test_hexagon/infrastructure.py | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/tests/python/contrib/test_hexagon/infrastructure.py b/tests/python/contrib/test_hexagon/infrastructure.py index 657290391a86..a1fbfdefcdbd 100644 --- a/tests/python/contrib/test_hexagon/infrastructure.py +++ b/tests/python/contrib/test_hexagon/infrastructure.py @@ -234,18 +234,20 @@ def compute(n, ho, wo, ko, hi, wi, ki): 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": - n, h, w, c = arr_np.shape 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 ["nhwc-4h2w32c2w-2d"]: + n, h, w, c = arr_np.shape return arr_np.reshape([n, h // 4, 4, 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]) if new_layout == "nc-1024-2d":