From 65638d2a6fd078891725eb6393fd073e3e6a8249 Mon Sep 17 00:00:00 2001 From: Chris Sullivan Date: Wed, 9 Nov 2022 13:57:46 -0800 Subject: [PATCH] Make 'allocate_hexagon_array' a hexagon contrib API --- python/tvm/contrib/hexagon/tools.py | 39 +++++++++++++++++++ .../contrib/test_hexagon/infrastructure.py | 38 ------------------ .../test_hexagon/test_2d_physical_buffers.py | 3 +- .../test_hexagon/test_benchmark_maxpool2d.py | 4 +- .../contrib/test_hexagon/test_memory_alloc.py | 3 +- .../contrib/test_hexagon/test_sigmoid.py | 3 +- .../topi/slice_op/test_argmax_slice.py | 4 +- .../topi/slice_op/test_avg_pool2d_slice.py | 2 +- .../topi/slice_op/test_cast_slice.py | 3 +- .../topi/slice_op/test_clip_slice.py | 3 +- .../topi/slice_op/test_conv2d_slice.py | 3 +- .../slice_op/test_depthwise_conv2d_slice.py | 3 +- .../topi/slice_op/test_dequantize_slice.py | 2 +- .../topi/slice_op/test_max_pool2d_slice.py | 3 +- .../topi/slice_op/test_relu_slice.py | 3 +- .../topi/slice_op/test_softmax_slice.py | 3 +- .../topi/slice_op/test_tanh_slice.py | 4 +- .../topi/test_add_subtract_multiply.py | 2 +- .../test_hexagon/topi/test_depth_to_space.py | 3 +- .../test_hexagon/topi/test_quantize.py | 2 +- .../contrib/test_hexagon/topi/test_reshape.py | 3 +- .../test_hexagon/topi/test_resize2d.py | 4 +- 22 files changed, 77 insertions(+), 60 deletions(-) mode change 100644 => 100755 tests/python/contrib/test_hexagon/test_2d_physical_buffers.py diff --git a/python/tvm/contrib/hexagon/tools.py b/python/tvm/contrib/hexagon/tools.py index 8c37261744d5..1c6468a0f5c7 100644 --- a/python/tvm/contrib/hexagon/tools.py +++ b/python/tvm/contrib/hexagon/tools.py @@ -20,6 +20,7 @@ import os import pathlib from typing import Union +import numpy import tvm import tvm.contrib.cc as cc @@ -203,3 +204,41 @@ def export_module(module, out_dir, binary_name="test_binary.so"): binary_path = pathlib.Path(out_dir) / binary_name module.save(str(binary_path)) return binary_path + + +def allocate_hexagon_array( + dev, tensor_shape=None, dtype=None, data=None, axis_separators=None, mem_scope=None +): + """ + Allocate a hexagon array which could be a 2D array + on physical memory defined by axis_separators + """ + if tensor_shape is None: + assert data is not None, "Must provide either tensor shape or numpy data array" + tensor_shape = data.shape + elif data is not None: + assert ( + tensor_shape == data.shape + ), "Mismatch between provided tensor shape and numpy data array shape" + + if dtype is None: + assert data is not None, "Must provide either dtype or numpy data array" + dtype = data.dtype.name + elif data is not None: + assert dtype == data.dtype, "Mismatch between provided dtype and numpy data array dtype" + + if axis_separators is None: + axis_separators = [] + + boundaries = [0, *axis_separators, len(tensor_shape)] + physical_shape = [ + 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: + arr.copyfrom(data.reshape(physical_shape)) + + return arr._create_view(tensor_shape) diff --git a/tests/python/contrib/test_hexagon/infrastructure.py b/tests/python/contrib/test_hexagon/infrastructure.py index 6f7e1904da2f..c04631156f1e 100644 --- a/tests/python/contrib/test_hexagon/infrastructure.py +++ b/tests/python/contrib/test_hexagon/infrastructure.py @@ -23,44 +23,6 @@ from tvm import te -def allocate_hexagon_array( - dev, tensor_shape=None, dtype=None, data=None, axis_separators=None, mem_scope=None -): - """ - Allocate a hexagon array which could be a 2D array - on physical memory defined by axis_separators - """ - if tensor_shape is None: - assert data is not None, "Must provide either tensor shape or numpy data array" - tensor_shape = data.shape - elif data is not None: - assert ( - tensor_shape == data.shape - ), "Mismatch between provided tensor shape and numpy data array shape" - - if dtype is None: - assert data is not None, "Must provide either dtype or numpy data array" - dtype = data.dtype.name - elif data is not None: - assert dtype == data.dtype, "Mismatch between provided dtype and numpy data array dtype" - - if axis_separators is None: - axis_separators = [] - - boundaries = [0, *axis_separators, len(tensor_shape)] - physical_shape = [ - 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: - arr.copyfrom(data.reshape(physical_shape)) - - return arr._create_view(tensor_shape) - - def ceildiv(o, d): assert o >= 0 assert d >= 0 diff --git a/tests/python/contrib/test_hexagon/test_2d_physical_buffers.py b/tests/python/contrib/test_hexagon/test_2d_physical_buffers.py old mode 100644 new mode 100755 index 7804ae2e4898..fb41e99a9bcb --- a/tests/python/contrib/test_hexagon/test_2d_physical_buffers.py +++ b/tests/python/contrib/test_hexagon/test_2d_physical_buffers.py @@ -31,8 +31,9 @@ from tvm import te from tvm.contrib.hexagon.pytest_plugin import requires_hexagon_toolchain from tvm.tir.stmt_functor import post_order_visit +from tvm.contrib.hexagon import allocate_hexagon_array -from .infrastructure import allocate_hexagon_array, get_hexagon_target +from .infrastructure import get_hexagon_target # Disabling invalid name as pylint assumes global variables as constants and # expects them to be all upper-case. Since these are used as diff --git a/tests/python/contrib/test_hexagon/test_benchmark_maxpool2d.py b/tests/python/contrib/test_hexagon/test_benchmark_maxpool2d.py index 24d1a3f788cf..42c77a9c9d2d 100644 --- a/tests/python/contrib/test_hexagon/test_benchmark_maxpool2d.py +++ b/tests/python/contrib/test_hexagon/test_benchmark_maxpool2d.py @@ -53,9 +53,9 @@ from tvm import te, topi, tir from tvm.topi import testing from tvm.contrib.hexagon.session import Session +from tvm.contrib.hexagon import allocate_hexagon_array - -from .infrastructure import allocate_hexagon_array, get_hexagon_target +from .infrastructure import get_hexagon_target from . import benchmark_util as bu # Pytest seems to require that fixture names exist in the current module. diff --git a/tests/python/contrib/test_hexagon/test_memory_alloc.py b/tests/python/contrib/test_hexagon/test_memory_alloc.py index f44e3cd0dc36..a0e3255a5428 100644 --- a/tests/python/contrib/test_hexagon/test_memory_alloc.py +++ b/tests/python/contrib/test_hexagon/test_memory_alloc.py @@ -20,8 +20,9 @@ import tvm from tvm.script import tir as T +from tvm.contrib.hexagon import allocate_hexagon_array -from .infrastructure import allocate_hexagon_array, get_hexagon_target +from .infrastructure import get_hexagon_target def generated_func(shape: tuple, dtype: str, axis_separators: list): diff --git a/tests/python/contrib/test_hexagon/test_sigmoid.py b/tests/python/contrib/test_hexagon/test_sigmoid.py index e115b188a3f0..cc633795c217 100644 --- a/tests/python/contrib/test_hexagon/test_sigmoid.py +++ b/tests/python/contrib/test_hexagon/test_sigmoid.py @@ -23,8 +23,9 @@ from tvm import te from tvm import tir from tvm import topi +from tvm.contrib.hexagon import allocate_hexagon_array -from .infrastructure import allocate_hexagon_array, get_hexagon_target +from .infrastructure import get_hexagon_target def sigmoid_compute(sigmoid_input): diff --git a/tests/python/contrib/test_hexagon/topi/slice_op/test_argmax_slice.py b/tests/python/contrib/test_hexagon/topi/slice_op/test_argmax_slice.py index 5f4a594fcfb1..92a951106765 100644 --- a/tests/python/contrib/test_hexagon/topi/slice_op/test_argmax_slice.py +++ b/tests/python/contrib/test_hexagon/topi/slice_op/test_argmax_slice.py @@ -22,7 +22,9 @@ from tvm import te import tvm.topi.hexagon.slice_ops as sl import tvm.contrib.hexagon -from ...infrastructure import allocate_hexagon_array, transform_numpy, get_hexagon_target +from tvm.contrib.hexagon import allocate_hexagon_array + +from ...infrastructure import transform_numpy, get_hexagon_target class TestArgMaxSlice: diff --git a/tests/python/contrib/test_hexagon/topi/slice_op/test_avg_pool2d_slice.py b/tests/python/contrib/test_hexagon/topi/slice_op/test_avg_pool2d_slice.py index 13876da87295..0eedfdbf8da1 100644 --- a/tests/python/contrib/test_hexagon/topi/slice_op/test_avg_pool2d_slice.py +++ b/tests/python/contrib/test_hexagon/topi/slice_op/test_avg_pool2d_slice.py @@ -23,8 +23,8 @@ from tvm.contrib.hexagon.session import Session import tvm.topi.hexagon.slice_ops as sl import tvm.topi.hexagon.qnn as qn +from tvm.contrib.hexagon import allocate_hexagon_array from ...infrastructure import ( - allocate_hexagon_array, transform_numpy, quantize_np, get_hexagon_target, diff --git a/tests/python/contrib/test_hexagon/topi/slice_op/test_cast_slice.py b/tests/python/contrib/test_hexagon/topi/slice_op/test_cast_slice.py index 3118c7be8efb..77776bc8da0b 100644 --- a/tests/python/contrib/test_hexagon/topi/slice_op/test_cast_slice.py +++ b/tests/python/contrib/test_hexagon/topi/slice_op/test_cast_slice.py @@ -22,8 +22,9 @@ import tvm.testing from tvm import te import tvm.topi.hexagon.slice_ops as sl +from tvm.contrib.hexagon import allocate_hexagon_array -from ...infrastructure import allocate_hexagon_array, transform_numpy, get_hexagon_target +from ...infrastructure import transform_numpy, get_hexagon_target class TestCastF16F32Slice2d: diff --git a/tests/python/contrib/test_hexagon/topi/slice_op/test_clip_slice.py b/tests/python/contrib/test_hexagon/topi/slice_op/test_clip_slice.py index e0a2e20a0b6b..d3f9804cd6c3 100644 --- a/tests/python/contrib/test_hexagon/topi/slice_op/test_clip_slice.py +++ b/tests/python/contrib/test_hexagon/topi/slice_op/test_clip_slice.py @@ -22,8 +22,9 @@ from tvm import te import tvm.testing import tvm.topi.hexagon.slice_ops as sl +from tvm.contrib.hexagon import allocate_hexagon_array -from ...infrastructure import allocate_hexagon_array, transform_numpy, get_hexagon_target +from ...infrastructure import transform_numpy, get_hexagon_target input_layout = tvm.testing.parameter( "nhwc-8h2w32c2w-2d", diff --git a/tests/python/contrib/test_hexagon/topi/slice_op/test_conv2d_slice.py b/tests/python/contrib/test_hexagon/topi/slice_op/test_conv2d_slice.py index c314e9655c9a..e06636cde365 100644 --- a/tests/python/contrib/test_hexagon/topi/slice_op/test_conv2d_slice.py +++ b/tests/python/contrib/test_hexagon/topi/slice_op/test_conv2d_slice.py @@ -24,8 +24,9 @@ import tvm.testing from tvm.topi.hexagon.slice_ops.conv2d import conv2d_compute, conv2d_schedule from tvm.topi.testing import conv2d_nhwc_python +from tvm.contrib.hexagon import allocate_hexagon_array -from ...infrastructure import allocate_hexagon_array, transform_numpy, get_hexagon_target +from ...infrastructure import transform_numpy, get_hexagon_target input_layout = tvm.testing.parameter( "nhwc-8h2w32c2w-2d", diff --git a/tests/python/contrib/test_hexagon/topi/slice_op/test_depthwise_conv2d_slice.py b/tests/python/contrib/test_hexagon/topi/slice_op/test_depthwise_conv2d_slice.py index 74e4d05446ed..e5a22e8879b5 100644 --- a/tests/python/contrib/test_hexagon/topi/slice_op/test_depthwise_conv2d_slice.py +++ b/tests/python/contrib/test_hexagon/topi/slice_op/test_depthwise_conv2d_slice.py @@ -25,8 +25,9 @@ import tvm.topi.hexagon.qnn as qn from tvm.topi.testing import depthwise_conv2d_python_nhwc from tvm.topi.hexagon.slice_ops.dwconv2d import dwconv2d_compute, dwconv2d_schedule +from tvm.contrib.hexagon import allocate_hexagon_array -from ...infrastructure import allocate_hexagon_array, transform_numpy, quantize_np +from ...infrastructure import transform_numpy, quantize_np @tvm.testing.fixture diff --git a/tests/python/contrib/test_hexagon/topi/slice_op/test_dequantize_slice.py b/tests/python/contrib/test_hexagon/topi/slice_op/test_dequantize_slice.py index 9b1c5bc5f614..8b9f49458df2 100644 --- a/tests/python/contrib/test_hexagon/topi/slice_op/test_dequantize_slice.py +++ b/tests/python/contrib/test_hexagon/topi/slice_op/test_dequantize_slice.py @@ -23,8 +23,8 @@ import tvm.testing from tvm import te from tvm.topi.hexagon import qnn +from tvm.contrib.hexagon import allocate_hexagon_array from ...infrastructure import ( - allocate_hexagon_array, transform_numpy, quantize_np, get_hexagon_target, diff --git a/tests/python/contrib/test_hexagon/topi/slice_op/test_max_pool2d_slice.py b/tests/python/contrib/test_hexagon/topi/slice_op/test_max_pool2d_slice.py index fcb4411609b2..4cd92f4dd27d 100644 --- a/tests/python/contrib/test_hexagon/topi/slice_op/test_max_pool2d_slice.py +++ b/tests/python/contrib/test_hexagon/topi/slice_op/test_max_pool2d_slice.py @@ -22,8 +22,9 @@ import tvm.testing from tvm.contrib.hexagon.session import Session import tvm.topi.hexagon.slice_ops as sl +from tvm.contrib.hexagon import allocate_hexagon_array -from ...infrastructure import allocate_hexagon_array, transform_numpy, get_hexagon_target +from ...infrastructure import transform_numpy, get_hexagon_target from ...pytest_util import ( get_multitest_ids, create_populated_numpy_ndarray, diff --git a/tests/python/contrib/test_hexagon/topi/slice_op/test_relu_slice.py b/tests/python/contrib/test_hexagon/topi/slice_op/test_relu_slice.py index 93a8d77827bf..1430551df719 100644 --- a/tests/python/contrib/test_hexagon/topi/slice_op/test_relu_slice.py +++ b/tests/python/contrib/test_hexagon/topi/slice_op/test_relu_slice.py @@ -21,8 +21,9 @@ import tvm.testing from tvm.topi.hexagon.slice_ops.relu import relu_compute, relu_stir_schedule from tvm import te +from tvm.contrib.hexagon import allocate_hexagon_array -from ...infrastructure import allocate_hexagon_array, transform_numpy, get_hexagon_target +from ...infrastructure import transform_numpy, get_hexagon_target @tvm.testing.fixture diff --git a/tests/python/contrib/test_hexagon/topi/slice_op/test_softmax_slice.py b/tests/python/contrib/test_hexagon/topi/slice_op/test_softmax_slice.py index a3db1b6dcdbe..2707ed3a5af1 100644 --- a/tests/python/contrib/test_hexagon/topi/slice_op/test_softmax_slice.py +++ b/tests/python/contrib/test_hexagon/topi/slice_op/test_softmax_slice.py @@ -20,8 +20,7 @@ from tvm import te from tvm.topi.testing import softmax_python import tvm.topi.hexagon.slice_ops as sl - -from ...infrastructure import allocate_hexagon_array +from tvm.contrib.hexagon import allocate_hexagon_array def transform_numpy(arr_np, layout): diff --git a/tests/python/contrib/test_hexagon/topi/slice_op/test_tanh_slice.py b/tests/python/contrib/test_hexagon/topi/slice_op/test_tanh_slice.py index f8c14ef934a1..6297ef2c1e6e 100644 --- a/tests/python/contrib/test_hexagon/topi/slice_op/test_tanh_slice.py +++ b/tests/python/contrib/test_hexagon/topi/slice_op/test_tanh_slice.py @@ -22,7 +22,9 @@ from tvm import te import tvm.topi.hexagon.slice_ops as sl import tvm.contrib.hexagon -from ...infrastructure import allocate_hexagon_array, transform_numpy, get_hexagon_target +from tvm.contrib.hexagon import allocate_hexagon_array + +from ...infrastructure import transform_numpy, get_hexagon_target # pylint: disable=invalid-name diff --git a/tests/python/contrib/test_hexagon/topi/test_add_subtract_multiply.py b/tests/python/contrib/test_hexagon/topi/test_add_subtract_multiply.py index d689888d6e85..e0bb6b5864d3 100644 --- a/tests/python/contrib/test_hexagon/topi/test_add_subtract_multiply.py +++ b/tests/python/contrib/test_hexagon/topi/test_add_subtract_multiply.py @@ -21,8 +21,8 @@ from tvm import te import tvm.topi.hexagon.slice_ops as sl import tvm.topi.hexagon.qnn as qn +from tvm.contrib.hexagon import allocate_hexagon_array from ..infrastructure import ( - allocate_hexagon_array, transform_numpy, quantize_np, get_hexagon_target, diff --git a/tests/python/contrib/test_hexagon/topi/test_depth_to_space.py b/tests/python/contrib/test_hexagon/topi/test_depth_to_space.py index 0cb41b595255..7d4afb953a50 100644 --- a/tests/python/contrib/test_hexagon/topi/test_depth_to_space.py +++ b/tests/python/contrib/test_hexagon/topi/test_depth_to_space.py @@ -25,8 +25,9 @@ import tvm.testing from tvm.topi.hexagon.slice_ops.depth_to_space import d2s_compute, d2s_schedule from tvm.topi.testing import depth_to_space_python +from tvm.contrib.hexagon import allocate_hexagon_array -from ..infrastructure import allocate_hexagon_array, transform_numpy, get_hexagon_target +from ..infrastructure import transform_numpy, get_hexagon_target class TestD2SSlice: diff --git a/tests/python/contrib/test_hexagon/topi/test_quantize.py b/tests/python/contrib/test_hexagon/topi/test_quantize.py index a188f7cb2fe1..ac4f4d4e3047 100644 --- a/tests/python/contrib/test_hexagon/topi/test_quantize.py +++ b/tests/python/contrib/test_hexagon/topi/test_quantize.py @@ -20,8 +20,8 @@ import tvm from tvm import te import tvm.topi.hexagon.qnn as s1 +from tvm.contrib.hexagon import allocate_hexagon_array from ..infrastructure import ( - allocate_hexagon_array, transform_numpy, quantize_np, get_hexagon_target, diff --git a/tests/python/contrib/test_hexagon/topi/test_reshape.py b/tests/python/contrib/test_hexagon/topi/test_reshape.py index 33bb31902eaa..51ac12506023 100644 --- a/tests/python/contrib/test_hexagon/topi/test_reshape.py +++ b/tests/python/contrib/test_hexagon/topi/test_reshape.py @@ -21,8 +21,9 @@ import tvm.testing import tvm.topi.hexagon.slice_ops as sl from tvm import te +from tvm.contrib.hexagon import allocate_hexagon_array -from ..infrastructure import allocate_hexagon_array, transform_numpy, get_hexagon_target +from ..infrastructure import transform_numpy, get_hexagon_target BATCH_FLATTEN_FP16_TESTS = ( ([1, 1, 1, 2048], [1, 2048], "nhwc-1024c-2d", "nc-1024-2d", "float16"), diff --git a/tests/python/contrib/test_hexagon/topi/test_resize2d.py b/tests/python/contrib/test_hexagon/topi/test_resize2d.py index 44d9c95a2f06..c0c6e7ca0fb4 100644 --- a/tests/python/contrib/test_hexagon/topi/test_resize2d.py +++ b/tests/python/contrib/test_hexagon/topi/test_resize2d.py @@ -22,7 +22,9 @@ from tvm import te from tvm.topi.testing import resize2d_python import tvm.topi.hexagon as s1 -from ..infrastructure import allocate_hexagon_array, transform_numpy, get_hexagon_target +from tvm.contrib.hexagon import allocate_hexagon_array + +from ..infrastructure import transform_numpy, get_hexagon_target class TestResize2d: