[Hexagon] Lint tests part 2

tests/lint/pylint.sh

@@ -29,7 +29,8 @@ python3 -m pylint tests/python/integration/ --rcfile="$(dirname "$0")"/pylintrc
 # tests/python/contrib/test_hexagon tests
 python3 -m pylint tests/python/contrib/test_hexagon/*.py --rcfile="$(dirname "$0")"/pylintrc
 python3 -m pylint tests/python/contrib/test_hexagon/conv2d/*.py --rcfile="$(dirname "$0")"/pylintrc
-
+python3 -m pylint tests/python/contrib/test_hexagon/topi/*.py --rcfile="$(dirname "$0")"/pylintrc
+python3 -m pylint tests/python/contrib/test_hexagon/metaschedule_e2e/*.py --rcfile="$(dirname "$0")"/pylintrc
 
 # tests/python/frontend tests
 python3 -m pylint tests/python/frontend/caffe/test_forward.py --rcfile="$(dirname "$0")"/pylintrc

tests/python/contrib/test_hexagon/metaschedule_e2e/export_models.py

@@ -14,6 +14,8 @@
 # KIND, either express or implied.  See the License for the
 # specific language governing permissions and limitations
 # under the License.
+"""Hexagon MetaSchedule test helper functions."""
+
 import torch
 from torchvision.models import resnet
 from torchvision.models.quantization import resnet as qresnet
@@ -23,6 +25,7 @@
 
 
 def export_resnet50_fp16():
+    """Export Resnet50 FP16."""
     model = resnet.resnet50(pretrained=True).eval()
 
     pt_inp = torch.randn(1, 3, 224, 224)
@@ -34,14 +37,16 @@ def export_resnet50_fp16():
     mod, params = relay.frontend.from_pytorch(script_module, input_shapes)
     mod = relay.transform.ToMixedPrecision("float16")(mod)
 
-    with open("resnet50_fp16.json", "w") as fo:
-        fo.write(tvm.ir.save_json(mod))
+    with open("resnet50_fp16.json", "w") as file:
+        file.write(tvm.ir.save_json(mod))
 
-    with open("resnet50_fp16.params", "wb") as fo:
-        fo.write(relay.save_param_dict(params))
+    with open("resnet50_fp16.params", "wb") as file:
+        file.write(relay.save_param_dict(params))
 
 
 def export_resnet50_int8():
+    """Export Resnet50 INT8."""
+
     def quantize_model(model, inp):
         model.fuse_model()
         model.qconfig = torch.quantization.get_default_qconfig("fbgemm")
@@ -62,11 +67,11 @@ def quantize_model(model, inp):
         script_module, input_shapes, keep_quantized_weight=True
     )
 
-    with open("resnet50_int8.json", "w") as fo:
-        fo.write(tvm.ir.save_json(mod))
+    with open("resnet50_int8.json", "w") as file:
+        file.write(tvm.ir.save_json(mod))
 
-    with open("resnet50_int8.params", "wb") as fo:
-        fo.write(relay.save_param_dict(params))
+    with open("resnet50_int8.params", "wb") as file:
+        file.write(relay.save_param_dict(params))
 
 
 if __name__ == "__main__":

tests/python/contrib/test_hexagon/metaschedule_e2e/test_resnet50_fp16.py

@@ -14,10 +14,12 @@
 # KIND, either express or implied.  See the License for the
 # specific language governing permissions and limitations
 # under the License.
+"""Test Resnet50 float16 with MetaSchedule"""
+
 import os
-import pytest
 import tempfile
 
+import pytest
 import numpy as np
 
 import tvm.testing
@@ -29,12 +31,6 @@
 from ..infrastructure import get_hexagon_target
 
 
-target = get_hexagon_target("v69")
-target_llvm = tvm.target.Target("llvm")
-model_json = "resnet50_fp16.json"
-model_params = "resnet50_fp16.params"
-
-
 def convert_conv2d_layout(mod, desired_layouts):
     with tvm.transform.PassContext(opt_level=3):
         seq = tvm.transform.Sequential([relay.transform.ConvertLayout(desired_layouts)])
@@ -44,14 +40,20 @@ def convert_conv2d_layout(mod, desired_layouts):
 @pytest.mark.skip("End-to-end tuning is skipped on CI.")
 @tvm.testing.requires_hexagon
 def test_resnet50(hexagon_launcher):
+    """Test Resnet50."""
+    model_json = "resnet50_fp16.json"
+    target_llvm = tvm.target.Target("llvm")
+    target_hexagon = get_hexagon_target("v69")
+    model_params = "resnet50_fp16.params"
+
     if not os.path.exists(model_json):
         pytest.skip(msg="Run python export_models.py first.")
 
-    with open(model_json, "r") as fi:
-        mod = tvm.ir.load_json(fi.read())
+    with open(model_json, "r") as file:
+        mod = tvm.ir.load_json(file.read())
 
-    with open(model_params, "rb") as fi:
-        params = relay.load_param_dict(fi.read())
+    with open(model_params, "rb") as file:
+        params = relay.load_param_dict(file.read())
 
     mod = convert_conv2d_layout(mod, {"nn.conv2d": ["NHWC", "HWIO"]})
 
@@ -66,7 +68,7 @@ def test_resnet50(hexagon_launcher):
     with tempfile.TemporaryDirectory() as work_dir:
         database = ms.relay_integration.tune_relay(
             mod=mod,
-            target=target,
+            target=target_hexagon,
             params=params,
             work_dir=work_dir,
             # for faster tuning
@@ -88,7 +90,7 @@ def test_resnet50(hexagon_launcher):
         hexagon_lowered = ms.relay_integration.compile_relay(
             database=database,
             mod=mod,
-            target=target,
+            target=target_hexagon,
             params=params,
         )
 
@@ -127,3 +129,7 @@ def test_resnet50(hexagon_launcher):
             hexagon_lowered.get_graph_json(), hexagon_lowered.lib
         )
         print(debug_ex.profile(input_name=inp.copy()))
+
+
+if __name__ == "__main__":
+    tvm.testing.main()

tests/python/contrib/test_hexagon/metaschedule_e2e/test_resnet50_int8.py

@@ -14,12 +14,15 @@
 # KIND, either express or implied.  See the License for the
 # specific language governing permissions and limitations
 # under the License.
+"""Test Resnet50 int8 with MetaSchedule"""
+
 import os
-import numpy as np
-import pytest
 import tempfile
 from typing import Optional
 
+import numpy as np
+import pytest
+
 import tvm
 import tvm.testing
 from tvm import relay
@@ -31,15 +34,15 @@
 from tvm.tir.schedule import BlockRV, Schedule
 from ..infrastructure import get_hexagon_target
 
-
-executor = relay.backend.Executor("graph", {"link-params": True})
-target = get_hexagon_target("v68")
-target_llvm = tvm.target.Target("llvm")
-model_json = "resnet50_int8.json"
-model_params = "resnet50_int8.params"
+MODEL_JSON = "resnet50_int8.json"
+EXECUTOR = relay.backend.Executor("graph", {"link-params": True})
+TARGET_LLVM = tvm.target.Target("llvm")
+TARGET_HEXAGON = get_hexagon_target("v68")
+MODEL_PARAMS = "resnet50_int8.params"
 
 
 def tune_vrmpy_auto_tensorize(mod, params, hexagon_launcher):
+    """Tune VRMPY with auto tensorization."""
     sch_rules = [
         schedule_rule.AutoInline(
             into_producer=False,
@@ -95,12 +98,12 @@ def tune_vrmpy_auto_tensorize(mod, params, hexagon_launcher):
 
     # This line is necessary for link-params to take effect during
     # task extraction and relay.build(...).
-    mod = mod.with_attr("executor", executor)
+    mod = mod.with_attr("executor", EXECUTOR)
 
     with tempfile.TemporaryDirectory() as work_dir:
         database = ms.relay_integration.tune_relay(
             mod=mod,
-            target=target,
+            target=TARGET_HEXAGON,
             params=params,
             work_dir=work_dir,
             # for faster tuning
@@ -129,22 +132,23 @@ def tune_vrmpy_auto_tensorize(mod, params, hexagon_launcher):
         return ms.relay_integration.compile_relay(
             database=database,
             mod=mod,
-            target=target,
+            target=TARGET_HEXAGON,
             params=params,
         )
 
 
 @pytest.mark.skip("End-to-end tuning is skipped on CI.")
 @tvm.testing.requires_hexagon
 def test_resnet50(hexagon_launcher):
-    if not os.path.exists(model_json):
+    """Test Resnet50."""
+    if not os.path.exists(MODEL_JSON):
         pytest.skip(msg="Run python export_models.py first.")
 
-    with open(model_json, "r") as fi:
-        mod = tvm.ir.load_json(fi.read())
+    with open(MODEL_JSON, "r") as file:
+        mod = tvm.ir.load_json(file.read())
 
-    with open(model_params, "rb") as fi:
-        params = relay.load_param_dict(fi.read())
+    with open(MODEL_PARAMS, "rb") as file:
+        params = relay.load_param_dict(file.read())
     inp = np.random.randn(1, 3, 224, 224).astype("float32")
     input_name = "image"
 
@@ -156,15 +160,15 @@ def test_resnet50(hexagon_launcher):
         with tvm.transform.PassContext(opt_level=3):
             hexagon_lowered = relay.build(
                 mod,
-                tvm.target.Target(target, host=target),
+                tvm.target.Target(TARGET_HEXAGON, host=TARGET_HEXAGON),
                 params=params,
-                executor=executor,
+                executor=EXECUTOR,
             )
 
     with tvm.transform.PassContext(opt_level=3):
         llvm_lowered = tvm.relay.build(
             mod,
-            tvm.target.Target(target_llvm, host=target_llvm),
+            tvm.target.Target(TARGET_LLVM, host=TARGET_LLVM),
             params=params,
         )
 
@@ -191,16 +195,16 @@ def test_resnet50(hexagon_launcher):
         print(debug_ex.profile(input_name=inp.copy()))
 
 
-def _schedule_packed_8x8x32_conv2d(do_tune: bool):
+def _schedule_packed_8x8x32_conv2d():
     """Manually schedule a conv2d block, created from TE compute op via CreatePrimFunc,
     using 8x8x32 packed layout.
     """
 
     def schedule_fn(sch, conv2d_block: Optional[BlockRV] = None) -> bool:
-        if conv2d_block == None:
+        if conv2d_block is None:
             try:
                 conv2d_block = sch.get_block("conv2d_NCHWc_int8")
-            except:
+            except ValueError:
                 return False
 
         assert "conv2d_NCHWc_int8" in sch.get(conv2d_block).annotations["schedule_rule"]
@@ -234,13 +238,13 @@ def schedule_fn(sch, conv2d_block: Optional[BlockRV] = None) -> bool:
             # be desirable to do this with coarser spatial granularity
             sch.compute_at(conv2d_block, loops[4])
 
-        def index_map_nchw32c_nchw8h8w32c(n, c, h, w, c32):
-            return [n, c, h // 8, w // 8, h % 8, w % 8, c32]
+        def index_map_nchw32c_nchw8h8w32c(n_batch, channel, height, width, channel_32):
+            return [n_batch, channel, height // 8, width // 8, height % 8, width % 8, channel_32]
 
         # Add cache for input and output activation layout transform,
         # note that weight is already in correct layout
-        input_cache = sch.cache_read(conv2d_block, 0, "global")
-        output_cache = sch.cache_write(outer_block, 0, "global")
+        input_cache = sch.cache_read(conv2d_block, 0, "global")  # pylint: disable=unused-variable
+        output_cache = sch.cache_write(outer_block, 0, "global")  # pylint: disable=unused-variable
         # Transform the layout of the input
         sch.transform_layout(
             conv2d_block, ("read", 0), index_map=index_map_nchw32c_nchw8h8w32c, pad_value=0
@@ -259,23 +263,25 @@ def index_map_nchw32c_nchw8h8w32c(n, c, h, w, c32):
 
 
 def tune_packed_8x8x32_template(mod, params, hexagon_launcher):
+    """Generate packed 8*8*32 template."""
+
     def schedule_rule_conv2d_packed_8x8x32(sch: Schedule, conv2d_block: BlockRV):
-        _schedule_packed_8x8x32_conv2d(do_tune=True)(sch, conv2d_block)
+        _schedule_packed_8x8x32_conv2d()(sch, conv2d_block)
         return [sch]
 
     register_func("meta_schedule.conv2d_NCHWc_int8", schedule_rule_conv2d_packed_8x8x32)
 
     def schedule_conv2d_for_tune(sch: Schedule):
-        _schedule_packed_8x8x32_conv2d(do_tune=True)(sch)
+        _schedule_packed_8x8x32_conv2d()(sch)
 
     # This line is necessary for link-params to take effect during
     # task extraction and relay.build(...).
-    mod = mod.with_attr("executor", executor)
+    mod = mod.with_attr("executor", EXECUTOR)
 
     with tempfile.TemporaryDirectory() as work_dir:
         database = ms.relay_integration.tune_relay(
             mod=mod,
-            target=target,
+            target=TARGET_HEXAGON,
             params=params,
             work_dir=work_dir,
             max_trials_global=20000,
@@ -309,22 +315,23 @@ def schedule_conv2d_for_tune(sch: Schedule):
         return ms.relay_integration.compile_relay(
             database=database,
             mod=mod,
-            target=target,
+            target=TARGET_HEXAGON,
             params=params,
         )
 
 
 @pytest.mark.skip("End-to-end tuning is skipped on CI.")
 @tvm.testing.requires_hexagon
 def test_packed_8x8x32_resnet50(hexagon_launcher):
-    if not os.path.exists(model_json):
+    """Test packed 8*8*32 Resnet50"""
+    if not os.path.exists(MODEL_JSON):
         pytest.skip(msg="Run python export_models.py first.")
 
-    with open(model_json, "r") as fi:
-        mod = tvm.ir.load_json(fi.read())
+    with open(MODEL_JSON, "r") as file:
+        mod = tvm.ir.load_json(file.read())
 
-    with open(model_params, "rb") as fi:
-        params = relay.load_param_dict(fi.read())
+    with open(MODEL_PARAMS, "rb") as file:
+        params = relay.load_param_dict(file.read())
     inp = np.random.randn(1, 3, 224, 224).astype("float32")
     input_name = "image"
 
@@ -336,15 +343,15 @@ def test_packed_8x8x32_resnet50(hexagon_launcher):
         with tvm.transform.PassContext(opt_level=3):
             hexagon_lowered = relay.build(
                 mod,
-                tvm.target.Target(target, host=target),
+                tvm.target.Target(TARGET_HEXAGON, host=TARGET_HEXAGON),
                 params=params,
-                executor=executor,
+                executor=EXECUTOR,
             )
 
     with tvm.transform.PassContext(opt_level=3):
         llvm_lowered = tvm.relay.build(
             mod,
-            tvm.target.Target(target_llvm, host=target_llvm),
+            tvm.target.Target(TARGET_LLVM, host=TARGET_LLVM),
             params=params,
         )
 
@@ -360,3 +367,7 @@ def test_packed_8x8x32_resnet50(hexagon_launcher):
         ref_result = llvm_graph_mod.get_output(0).numpy()
 
         np.testing.assert_allclose(ref_result, hexagon_output, atol=1e-4, rtol=1e-5)
+
+
+if __name__ == "__main__":
+    tvm.testing.main()

tests/python/contrib/test_hexagon/topi/__init__.py

@@ -15,4 +15,4 @@
 # specific language governing permissions and limitations
 # under the License.
 
-""" Testing infrastructure for Hexagon/TOPI """
+""" Hexagon TOPI tests """

tests/python/contrib/test_hexagon/topi/slice_op/__init__.py

@@ -0,0 +1,18 @@
+# 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 TOPI Slice OP tests """