This is an automated email from the ASF dual-hosted git repository.
mbaret 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 d9fac4f [microNPU] adding more tests with USMP (#10362)
d9fac4f is described below
commit d9fac4f305c5462de4ab08cd220d60b42c0414f7
Author: Manupa Karunaratne <[email protected]>
AuthorDate: Fri Feb 25 11:28:50 2022 +0000
[microNPU] adding more tests with USMP (#10362)
Adding a few tests to confirm memory usage
with and without USMP.
- Supporting the toggle to disable storage_rewrite.
- There is a slight change to tir_to_cs_translator to
add index of Load nodes associated with NpuAddressRange objects
---
.../relay/backend/contrib/ethosu/tir/compiler.py | 4 +-
.../backend/contrib/ethosu/tir_to_cs_translator.py | 5 +-
tests/python/contrib/test_ethosu/infra.py | 2 +
tests/python/contrib/test_ethosu/test_networks.py | 81 ++++++++++++++--------
4 files changed, 62 insertions(+), 30 deletions(-)
diff --git a/python/tvm/relay/backend/contrib/ethosu/tir/compiler.py
b/python/tvm/relay/backend/contrib/ethosu/tir/compiler.py
index ee35da4..bdc3b31 100644
--- a/python/tvm/relay/backend/contrib/ethosu/tir/compiler.py
+++ b/python/tvm/relay/backend/contrib/ethosu/tir/compiler.py
@@ -87,7 +87,9 @@ def lower_ethosu(sch, args, const_dict, name="main"):
mod = ethosu_passes.ReplaceOperators()(mod)
mod = tvm.tir.transform.RemoveNoOp()(mod)
mod, const_dict = ethosu_passes.EncodeConstants(const_dict)(mod)
- mod = tvm.tir.transform.StorageRewrite()(mod)
+ disable_storage_rewrite = curr_cfg.get("tir.disable_storage_rewrite",
False)
+ if not disable_storage_rewrite:
+ mod = tvm.tir.transform.StorageRewrite()(mod)
mod = tvm.tir.transform.RemoveNoOp()(mod)
mod = ethosu_passes.AnnotateAllocates()(mod)
mod, const_dict =
ethosu_passes.CreatePrimFuncWithoutConstants(const_dict)(mod)
diff --git a/python/tvm/relay/backend/contrib/ethosu/tir_to_cs_translator.py
b/python/tvm/relay/backend/contrib/ethosu/tir_to_cs_translator.py
index c53c328..f642f5f 100644
--- a/python/tvm/relay/backend/contrib/ethosu/tir_to_cs_translator.py
+++ b/python/tvm/relay/backend/contrib/ethosu/tir_to_cs_translator.py
@@ -423,10 +423,13 @@ def assign_addresses(buffer_info, npu_ops,
scratch_region_map):
def replace_npu_address_range_with_address(npu_addr_range):
assert isinstance(npu_addr_range.address, tvm.tir.Load)
buffer = npu_addr_range.address.buffer_var
+ index = int(
+ npu_addr_range.address.index *
(np.iinfo(np.dtype(npu_addr_range.address)).bits // 8)
+ )
if buffer in scratch_region_map.keys():
return vapi.NpuAddressRange(
scratch_region_map[buffer].region,
- scratch_region_map[buffer].offset,
+ scratch_region_map[buffer].offset + index,
npu_addr_range.length,
)
assert buffer in buffer_addresses.keys(), f"searching for buffer :
{buffer}, but not found"
diff --git a/tests/python/contrib/test_ethosu/infra.py
b/tests/python/contrib/test_ethosu/infra.py
index b355c44..be5466c 100644
--- a/tests/python/contrib/test_ethosu/infra.py
+++ b/tests/python/contrib/test_ethosu/infra.py
@@ -216,6 +216,8 @@ def create_test_runner(accel="ethos-u55-256",
enable_usmp=True):
"accelerator_config": accel,
},
"tir.usmp.enable": enable_usmp,
+ "tir.usmp.algorithm": "hill_climb",
+ "tir.disable_storage_rewrite": enable_usmp,
},
)
diff --git a/tests/python/contrib/test_ethosu/test_networks.py
b/tests/python/contrib/test_ethosu/test_networks.py
index 7e3140f..b73ba4f 100644
--- a/tests/python/contrib/test_ethosu/test_networks.py
+++ b/tests/python/contrib/test_ethosu/test_networks.py
@@ -29,50 +29,75 @@ import tvm.micro as micro
from tvm import relay
from tvm.relay.backend.contrib.ethosu import util
from tvm.relay.op.contrib.ethosu import partition_for_ethosu
+from tests.python.relay.aot.aot_test_utils import
create_relay_module_and_inputs_from_tflite_file
+from tvm.micro import model_library_format as mlf
import tvm.relay.testing.tf as tf_testing
from . import infra
+MOBILENET_V1_URL = (
+
"https://storage.googleapis.com/download.tensorflow.org/models/mobilenet_v1_2018_08_02/mobilenet_v1_1.0_224_quant.tgz";,
+ "mobilenet_v1_1.0_224_quant.tflite",
+)
+
+MOBILENET_V2_URL = (
+
"https://storage.googleapis.com/download.tensorflow.org/models/tflite_11_05_08/mobilenet_v2_1.0_224_quant.tgz";,
+ "mobilenet_v2_1.0_224_quant.tflite",
+)
+
+
@pytest.mark.parametrize(
- "accel_type, enable_usmp",
+ "accel_type, model_url, workspace_size, tolerance",
[
- ("ethos-u55-256", True),
- ("ethos-u55-256", False),
- ("ethos-u65-256", True),
- ("ethos-u65-256", False),
- ("ethos-u55-128", True),
- ("ethos-u55-64", True),
- ("ethos-u55-32", True),
+ ("ethos-u65-256", MOBILENET_V1_URL, 1423344, 10),
+ ("ethos-u65-256", MOBILENET_V2_URL, 2185584, 5),
+ ("ethos-u55-256", MOBILENET_V1_URL, 1423344, 10),
+ ("ethos-u55-256", MOBILENET_V2_URL, 2185584, 5),
+ ("ethos-u55-128", MOBILENET_V2_URL, 2185584, 5),
+ ("ethos-u55-64", MOBILENET_V2_URL, 2185584, 5),
+ ("ethos-u55-32", MOBILENET_V2_URL, 2185584, 5),
],
)
-def test_forward_mobilenet_v1(accel_type, enable_usmp):
- """Test the Mobilenet V1 TF Lite model."""
+def test_networks_without_usmp(accel_type, model_url, workspace_size,
tolerance):
np.random.seed(23)
- tflite_model_file = tf_testing.get_workload_official(
- "https://storage.googleapis.com/download.tensorflow.org/";
- "models/mobilenet_v1_2018_08_02/mobilenet_v1_1.0_224_quant.tgz",
- "mobilenet_v1_1.0_224_quant.tflite",
+ tflite_model_file = tf_testing.get_workload_official(model_url[0],
model_url[1])
+ mod, input_data, params =
create_relay_module_and_inputs_from_tflite_file(tflite_model_file)
+ output_data = generate_ref_data(mod, input_data, params)
+ mod = partition_for_ethosu(mod, params)
+ compiled_models = infra.build_source(
+ mod, input_data, output_data, accel_type, output_tolerance=tolerance,
enable_usmp=False
)
- with open(tflite_model_file, "rb") as f:
- tflite_model_buf = f.read()
- input_tensor = "input"
- input_dtype = "uint8"
- input_shape = (1, 224, 224, 3)
- in_min, in_max = util.get_range_for_dtype_str(input_dtype)
- input_data = np.random.randint(in_min, high=in_max, size=input_shape,
dtype=input_dtype)
+ mlf_memory_map = mlf._build_function_memory_map(
+ compiled_models[0].executor_factory.function_metadata
+ )
+ assert mlf_memory_map["main"][0]["workspace_size_bytes"] == workspace_size
+ infra.verify_source(compiled_models, accel_type, enable_usmp=False)
- relay_mod, params = convert_to_relay(tflite_model_buf)
- input_data = {input_tensor: input_data}
- output_data = generate_ref_data(relay_mod, input_data)
- mod = partition_for_ethosu(relay_mod, params)
[email protected](
+ "accel_type, model_url, workspace_size, tolerance",
+ [
+ ("ethos-u65-256", MOBILENET_V1_URL, 1205872, 10),
+ ("ethos-u55-256", MOBILENET_V2_URL, 1507152, 5),
+ ],
+)
+def test_networks_with_usmp(accel_type, model_url, workspace_size, tolerance):
+ np.random.seed(23)
+ tflite_model_file = tf_testing.get_workload_official(model_url[0],
model_url[1])
+ mod, input_data, params =
create_relay_module_and_inputs_from_tflite_file(tflite_model_file)
+ output_data = generate_ref_data(mod, input_data, params)
+ mod = partition_for_ethosu(mod, params)
compiled_models = infra.build_source(
- mod, input_data, output_data, accel_type, output_tolerance=10,
enable_usmp=enable_usmp
+ mod, input_data, output_data, accel_type, output_tolerance=tolerance,
enable_usmp=True
)
- infra.verify_source(compiled_models, accel_type, enable_usmp=enable_usmp)
+ allocated_pool_info = list(
+
dict(compiled_models[0].executor_factory.executor_codegen_metadata.pool_inputs).values()
+ )[0]
+ assert allocated_pool_info.allocated_size == workspace_size
+ infra.verify_source(compiled_models, accel_type, enable_usmp=True)
if __name__ == "__main__":
- test_forward_mobilenet_v1()
+ pytest.main([__file__])
