This is an automated email from the ASF dual-hosted git repository.
tqchen pushed a commit to branch main
in repository https://gitbox.apache.org/repos/asf/tvm.git
The following commit(s) were added to refs/heads/main by this push:
new 39cc48bef8 [Hexagon][Tests] Clean up stale hexagon tests (#19747)
39cc48bef8 is described below
commit 39cc48bef8caac37862f5bec1de17f5ac7b10b76
Author: Shushi Hong <[email protected]>
AuthorDate: Fri Jun 12 08:30:31 2026 -0400
[Hexagon][Tests] Clean up stale hexagon tests (#19747)
The test_hexagon directory is not enrolled in CI and has drifted out of
sync with the codebase:
- Delete test_fixed_point_conversion.py: it imports
tvm.topi.hexagon.utils, but the tvm.topi.hexagon module (including
get_fixed_point_value under test) no longer exists in the source tree.
- Delete test_2d_physical_buffers.py: its ir_module fixture was removed
in #17665 (Feb 2025), leaving all 36 tests erroring at setup with
"fixture 'ir_module' not found" ever since.
- Fix test_take.py: drop the stale exec_mode="compiled" argument that
tvm.compile no longer accepts. Both tests in the file pass again (llvm
host target, no Hexagon hardware required).
---
.../test_hexagon/test_2d_physical_buffers.py | 355 ---------------------
.../test_hexagon/test_fixed_point_conversion.py | 70 ----
tests/python/contrib/test_hexagon/test_take.py | 2 +-
3 files changed, 1 insertion(+), 426 deletions(-)
diff --git a/tests/python/contrib/test_hexagon/test_2d_physical_buffers.py
b/tests/python/contrib/test_hexagon/test_2d_physical_buffers.py
deleted file mode 100644
index 5ffd60626a..0000000000
--- a/tests/python/contrib/test_hexagon/test_2d_physical_buffers.py
+++ /dev/null
@@ -1,355 +0,0 @@
-#!/usr/bin/env python3
-
-# 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.
-
-"""Test 2d physical buffers"""
-
-import contextlib
-
-import numpy as np
-import pytest
-
-import tvm
-
-# Needed to register the link_shared packedfunc.
-import tvm.contrib.hexagon
-import tvm.testing
-from tvm import te
-from tvm.contrib.hexagon import allocate_hexagon_array
-from tvm.contrib.hexagon.pytest_plugin import requires_hexagon_toolchain
-from tvm.tirx.stmt_functor import post_order_visit
-
-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
-# tvm.testing.parameters, if they are made upper-case, the functions which take
-# them as arguments would also need to be upper-case, and pylint would complain
-# there as well
-# pylint: disable=invalid-name
-
-schedule_type = tvm.testing.parameter("TE", "TIR")
-
-dtype = tvm.testing.parameter("int8")
-batch_size = tvm.testing.parameter(
- 16,
- 2,
-)
-input_channels = tvm.testing.parameter(
- 32,
-)
-input_image_shape = tvm.testing.parameter(
- by_dict={
- "8x8": (8, 8),
- "32x32": (32, 32),
- }
-)
-
-input_layout = tvm.testing.parameter(
- "nhwc",
- "nchw-8h8w32c-1d",
-)
-output_layout = tvm.testing.parameter(
- "nhwc",
- "nchw-8h8w32c-1d",
-)
-working_layout, working_scope = tvm.testing.parameters(
- ("nhwc", "global"),
- ("nhwc", "global.vtcm"),
- ("nchw-8h8w32c-1d", "global"),
- ("nchw-8h8w32c-1d", "global.vtcm"),
- # 2-d memory may only occur in vtcm memory
- ("nchw-8h8w32c-2d", "global.vtcm"),
-)
-
-# pylint: enable=invalid-name
-
-
[email protected]
-def target_host():
- """Return tvm target.Target with host attached"""
- return get_hexagon_target("v68")
-
-
-# Disabling redefined-outer-name for the whole file as there isn't any easy
-# solution yet to refactor tvm.testing.fixture fixtures that avoid redefining
-# outer variable names
-# pylint: disable=redefined-outer-name
-
-
[email protected]
-def input_shape(batch_size, input_channels, input_image_shape):
- return [batch_size, *input_image_shape, input_channels]
-
-
-def transform_shape(shape, layout):
- if layout == "nhwc":
- return shape
- if layout in ["nchw-8h8w32c-1d", "nchw-8h8w32c-2d"]:
- batch, height, width, channel = shape
- return [batch, (channel + 31) // 32, (height + 7) // 8, (width + 7) //
8, 8, 8, 32]
- raise RuntimeError(f"Unexpected layout '{layout}'")
-
-
-def transform_numpy(arr_np, layout):
- if layout == "nhwc":
- return arr_np
- if layout in ["nchw-8h8w32c-1d", "nchw-8h8w32c-2d"]:
- batch, height, width, channel = arr_np.shape
- return arr_np.reshape([batch, height // 8, 8, width // 8, 8, channel
// 32, 32]).transpose(
- 0, 5, 1, 3, 2, 4, 6
- )
- raise RuntimeError(f"Unexpected layout '{layout}'")
-
-
[email protected]
-def transformed_input_shape(input_shape, input_layout):
- return transform_shape(input_shape, input_layout)
-
-
[email protected]
-def transformed_output_shape(output_shape, output_layout):
- return transform_shape(output_shape, output_layout)
-
-
[email protected]
-def input_np(input_shape, dtype):
- return (100 * np.random.uniform(size=input_shape)).astype(dtype)
-
-
[email protected]
-def transformed_input_np(input_np, input_layout):
- return transform_numpy(input_np, input_layout)
-
-
[email protected]
-def transformed_expected_output_np(expected_output_np, output_layout):
- return transform_numpy(expected_output_np, output_layout)
-
-
-def layout_transform_1d(batch, height, width, channel):
- return [
- batch,
- channel // 32,
- height // 8,
- width // 8,
- height % 8,
- width % 8,
- channel % 32,
- ]
-
-
-def layout_transform_2d(batch, height, width, channel):
- return [
- batch,
- channel // 32,
- height // 8,
- width // 8,
- te.AXIS_SEPARATOR,
- height % 8,
- width % 8,
- channel % 32,
- ]
-
-
-def extract_buffers(stmt):
- buffers = []
-
- def visitor(node):
- if isinstance(node, tvm.tirx.BufferLoad | tvm.tirx.BufferStore):
- buffers.append(node.buffer)
-
- post_order_visit(stmt, visitor)
- return buffers
-
-
-class TestElementWise:
- """TestElementWise"""
-
- @tvm.testing.fixture
- def expected_output_np(self, input_np):
- return 2 * input_np
-
- @tvm.testing.fixture
- def output_shape(self, input_shape):
- return input_shape
-
- @tvm.testing.fixture
- def schedule_args(
- self,
- schedule_type,
- input_shape,
- dtype,
- input_layout,
- output_layout,
- working_layout,
- working_scope,
- ):
- """Create and return the schedule and input args after applying layout
transform"""
- if schedule_type == "TIR":
- return self._tir_schedule_args(
- input_shape, dtype, input_layout, output_layout,
working_layout, working_scope
- )
-
- else:
- raise ValueError(f"Unknown schedule type: {schedule_type}")
-
- def _te_tensors(self, input_shape, dtype):
- input_tensor = te.placeholder(input_shape, dtype, name="Input")
- output_tensor = te.compute(
- shape=input_tensor.shape,
- fcompute=lambda *indices: (2 *
input_tensor[indices]).astype(dtype),
- name="Output",
- )
- return input_tensor, output_tensor
-
- def _tir_schedule_args(
- self, input_shape, dtype, input_layout, output_layout, working_layout,
working_scope
- ):
- tensors = self._te_tensors(input_shape, dtype)
-
- sch = tvm.s_tir.Schedule(te.create_prim_func(tensors))
-
- cache_read_block = sch.cache_read("Output", 0, working_scope)
- cache_write_block = sch.cache_write("Output", 0, working_scope)
-
- def apply_transform(block, buffer_name, layout):
- if layout == "nhwc":
- pass
- elif layout == "nchw-8h8w32c-1d":
- sch.transform_layout(block, buffer_name, layout_transform_1d)
- elif layout == "nchw-8h8w32c-2d":
- sch.transform_layout(block, buffer_name, layout_transform_2d)
- else:
- raise RuntimeError(f"Unexpected layout '{layout}'")
-
- apply_transform(cache_read_block, ("read", 0), input_layout)
- apply_transform(cache_read_block, ("write", 0), working_layout)
- apply_transform(cache_write_block, ("read", 0), working_layout)
- apply_transform(cache_write_block, ("write", 0), output_layout)
-
- return [sch.mod]
-
- @tvm.testing.fixture
- def uses_unsupported_physical_dimensions( # pylint: disable=invalid-name
- self, target_host, input_layout, working_layout, output_layout
- ):
- uses_2d_memory = "nchw-8h8w32c-2d" in [input_layout, working_layout,
output_layout]
- can_handle_2d_memory = target_host.kind.name == "hexagon"
-
- return uses_2d_memory and not can_handle_2d_memory
-
- def test_param_shapes(self, ir_module, transformed_input_shape,
transformed_output_shape):
- func = ir_module["main"]
- primfunc_input_shape, primfunc_output_shape = [
- list(func.buffer_map[param].shape) for param in func.params
- ]
- assert primfunc_input_shape == transformed_input_shape
- assert primfunc_output_shape == transformed_output_shape
-
- def test_cache_shape(self, ir_module, input_layout, working_layout,
output_layout):
- """Test function to check expected_physical_dimensions for cached
buffers"""
- func = ir_module["main"]
- for buffer in extract_buffers(func.body):
- buffer_layout = {
- "Input": input_layout,
- "Input.global": working_layout,
- "Output.global": working_layout,
- "Input.global.vtcm": working_layout,
- "Output.global.vtcm": working_layout,
- "Output": output_layout,
- }[buffer.name.replace("_", ".")]
-
- expected_physical_dimensions = {
- "nhwc": 1,
- "nchw-8h8w32c-1d": 1,
- "nchw-8h8w32c-2d": 2,
- }[buffer_layout]
-
- assert len(buffer.shape) == expected_physical_dimensions
-
- @requires_hexagon_toolchain
- def test_build(self, schedule_args, target_host, input_layout,
working_layout, output_layout):
- """Testing build success/failure
-
- * On Hexagon targets, build must succeed for both 1-d and 2-d memory.
- * On non-Hexagon targets, build must succeed 1-d memory.
- * On non-Hexagon targets, build must fail and report an error for 2-d
memory.
- """
- # contextlib.nullcontext wasn't added until python3.7, and the
- # CI currently runs on python3.6. Therefore, using ExitStack
- # to manage an optional context instead.
- stack = contextlib.ExitStack()
-
- with stack:
- is_hexagon = target_host.kind.name == "hexagon"
- uses_2d_memory = "nchw-8h8w32c-2d" in [input_layout,
working_layout, output_layout]
- if uses_2d_memory and not is_hexagon:
- stack.enter_context(pytest.raises(RuntimeError))
-
- tvm.compile(*schedule_args, target=target_host)
-
- @tvm.testing.fixture
- def runtime_module(self, schedule_args, target_host):
- if target_host.kind.name != "hexagon":
- pytest.skip("Only running on hexagon")
-
- return tvm.compile(*schedule_args, target=target_host)
-
- @tvm.testing.requires_hexagon
- def test_execute(
- self,
- runtime_module,
- transformed_input_np,
- transformed_expected_output_np,
- input_layout,
- output_layout,
- hexagon_session,
- ):
- """Test execution of computes with 2d physical buffers"""
- if input_layout == "nchw-8h8w32c-2d":
- input_axis_separators = [4]
- else:
- input_axis_separators = []
-
- if output_layout == "nchw-8h8w32c-2d":
- output_axis_separators = [4]
- else:
- output_axis_separators = []
-
- input_arr = allocate_hexagon_array(
- hexagon_session.device,
- data=transformed_input_np,
- axis_separators=input_axis_separators,
- )
- output_arr = allocate_hexagon_array(
- hexagon_session.device,
- data=np.zeros_like(transformed_expected_output_np),
- axis_separators=output_axis_separators,
- )
-
- mod = hexagon_session.load_module(runtime_module)
-
- mod(input_arr, output_arr)
- output_np = output_arr.numpy()
-
- np.testing.assert_array_equal(output_np,
transformed_expected_output_np)
-
-
-if __name__ == "__main__":
- tvm.testing.main()
diff --git a/tests/python/contrib/test_hexagon/test_fixed_point_conversion.py
b/tests/python/contrib/test_hexagon/test_fixed_point_conversion.py
deleted file mode 100644
index c1bc517be7..0000000000
--- a/tests/python/contrib/test_hexagon/test_fixed_point_conversion.py
+++ /dev/null
@@ -1,70 +0,0 @@
-# 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.
-# ruff: noqa: F821
-
-"""
-Test float to fixed-point conversion. We do it by constructing a numpy array
with the
-wide range of floating-point values. These values are converted into the
-fixed-point value using topi.hexagon.utils.get_fixed_point_value. Then, these
values are
-converted back into float using scale_factor provided by the function. These
converted
-floating point values are then compared against the original values and an
assertion is
-raised if they happened to be outside of the expected tolerance.
-"""
-
-import math
-import struct
-
-import numpy as np
-
-from tvm.topi.hexagon import utils
-
-
-class TestFixedPointConversion:
- """Fixed point conversation test class"""
-
- def test_fixed_point_conversion(self):
- """Test fixed point conversion"""
- # Construct array with wide range of values
- fp1 = np.random.uniform(0.00001, 0.0002, size=(10))
- fp2 = np.random.uniform(0.001, 0.02, size=(10))
- fp3 = np.random.uniform(1, 20, size=(10))
- fp4 = np.random.uniform(900, 1000, size=(10))
- fp5 = np.random.uniform(1e9, 1e10, size=(10))
-
- # Test for values with largest possible exponent as per IEEE-754
floating-point
- # standard (actual exp value = 127, stored exp value = 254).
- fp6 = np.random.uniform(2.4e38, 2.5e38, size=(1))
-
- # Test for very small floating-point values.
- fp7 = np.random.uniform(1.4e-34, 1.7e-34, size=(1))
-
- float_arr = np.concatenate((fp1, fp2, fp3, fp4, fp5, fp6, fp7))
- for flp in float_arr:
- fxp, rsh = utils.get_fixed_point_value(flp, "int16")
- # Compute scale_factor using rsh (rsh is log2 of the
scale_factor). While doing this,
- # we use IEEE-754 floating-point representation since rsh can be
negative or positive.
-
- scale = ((rsh + 127) & 0xFF) << 23 # Add bias (127) and position
it into exponent bits
- scale_i = struct.pack("I", scale) # Pack it as integer
- scale_f = struct.unpack("f", scale_i) # Unpack as float
-
- converted_flp = fxp / scale_f[0]
- assert math.isclose(flp, converted_flp, rel_tol=1e-2)
-
-
-if __name__ == "__main__":
- tvm.testing.main()
diff --git a/tests/python/contrib/test_hexagon/test_take.py
b/tests/python/contrib/test_hexagon/test_take.py
index 8acbf53126..9e46da1c17 100644
--- a/tests/python/contrib/test_hexagon/test_take.py
+++ b/tests/python/contrib/test_hexagon/test_take.py
@@ -370,7 +370,7 @@ def test_value():
after = generate_take_op.PassReplaceWithTakeOpPrimFuncs()(before)
target = tvm.target.Target("llvm", host="llvm")
- ex = tvm.compile(after, target, exec_mode="compiled")
+ ex = tvm.compile(after, target)
vm = relax.VirtualMachine(ex, tvm.cpu())
res = vm["main"](inp_quant)
