From a74d0fef35b0ddcd75f446eea83039a53cc38dd7 Mon Sep 17 00:00:00 2001 From: Lunderberg Date: Fri, 4 Jun 2021 10:34:25 -0700 Subject: [PATCH] [Codegen] Use "target.build.$TARGET_KIND" for all codegen functions. (#8071) * [Codegen] Use "target.build.$TARGET_KIND" for all codegen functions. - Removed special case for "micro_dev" target. Instead, register BuildCHost as both "target.build.c" and "target.build.micro_dev". - Renamed "target.build.build.aocl_sw_emu" to "target.build.aocl_sw_emu". Appears to be a typo introduced in #841725cc585 * [micro_dev] Removed references to non-existent micro_dev device_api.micro_dev was removed in 745e542e4deaf44f3d6e5665299aa85ef8f4a6b9, but several references still remained. Co-authored-by: Eric Lunderberg --- include/tvm/runtime/device_api.h | 2 - python/tvm/__init__.py | 2 +- python/tvm/_ffi/runtime_ctypes.py | 2 - python/tvm/autotvm/measure/measure_methods.py | 2 - python/tvm/runtime/__init__.py | 2 +- python/tvm/runtime/module.py | 3 - python/tvm/runtime/ndarray.py | 27 +- src/runtime/module.cc | 2 - src/target/codegen.cc | 8 +- src/target/source/codegen_aocl.cc | 2 +- tests/micro/test_runtime_micro_on_arm.py | 370 ------------------ 11 files changed, 9 insertions(+), 413 deletions(-) delete mode 100644 tests/micro/test_runtime_micro_on_arm.py diff --git a/include/tvm/runtime/device_api.h b/include/tvm/runtime/device_api.h index a493469a333d..58b9ff1932cc 100644 --- a/include/tvm/runtime/device_api.h +++ b/include/tvm/runtime/device_api.h @@ -257,8 +257,6 @@ inline const char* DeviceName(int type) { return "ext_dev"; case kDLWebGPU: return "webgpu"; - case kDLMicroDev: - return "micro_dev"; case kDLHexagon: return "hexagon"; default: diff --git a/python/tvm/__init__.py b/python/tvm/__init__.py index 77630730f03a..55a228882691 100644 --- a/python/tvm/__init__.py +++ b/python/tvm/__init__.py @@ -31,7 +31,7 @@ # tvm.runtime from .runtime.object import Object from .runtime.ndarray import device, cpu, cuda, gpu, opencl, cl, vulkan, metal, mtl -from .runtime.ndarray import vpi, rocm, ext_dev, micro_dev, hexagon +from .runtime.ndarray import vpi, rocm, ext_dev, hexagon from .runtime import ndarray as nd # tvm.error diff --git a/python/tvm/_ffi/runtime_ctypes.py b/python/tvm/_ffi/runtime_ctypes.py index efea47752f6d..450a356aebdf 100644 --- a/python/tvm/_ffi/runtime_ctypes.py +++ b/python/tvm/_ffi/runtime_ctypes.py @@ -173,7 +173,6 @@ class Device(ctypes.Structure): 9: "vpi", 10: "rocm", 12: "ext_dev", - 13: "micro_dev", 14: "hexagon", 15: "webgpu", } @@ -194,7 +193,6 @@ class Device(ctypes.Structure): "vpi": 9, "rocm": 10, "ext_dev": 12, - "micro_dev": 13, "hexagon": 14, "webgpu": 15, } diff --git a/python/tvm/autotvm/measure/measure_methods.py b/python/tvm/autotvm/measure/measure_methods.py index 60a26ecd7d81..f41795fb0810 100644 --- a/python/tvm/autotvm/measure/measure_methods.py +++ b/python/tvm/autotvm/measure/measure_methods.py @@ -276,8 +276,6 @@ def get_build_kwargs(self): if "cuda" in self.task.target.keys: kwargs["cuda_arch"] = "sm_" + "".join(dev.compute_version.split(".")) - if self.task.target.device_name == "micro_dev": - kwargs.setdefault("build_option", {})["tir.disable_vectorize"] = True return kwargs diff --git a/python/tvm/runtime/__init__.py b/python/tvm/runtime/__init__.py index 265dedb63b57..71563b508290 100644 --- a/python/tvm/runtime/__init__.py +++ b/python/tvm/runtime/__init__.py @@ -27,7 +27,7 @@ # function exposures from .object_generic import convert_to_object, convert, const from .ndarray import device, cpu, cuda, gpu, opencl, cl, vulkan, metal, mtl -from .ndarray import vpi, rocm, ext_dev, micro_dev +from .ndarray import vpi, rocm, ext_dev from .module import load_module, enabled, system_lib from .container import String from .params import save_param_dict, load_param_dict diff --git a/python/tvm/runtime/module.py b/python/tvm/runtime/module.py index f0f33e162559..8107ab5b87d2 100644 --- a/python/tvm/runtime/module.py +++ b/python/tvm/runtime/module.py @@ -470,9 +470,6 @@ def load_module(path, fmt=""): files = [tar_temp.relpath(x) for x in tar_temp.listdir()] _cc.create_shared(path + ".so", files, cc=cc) path += ".so" - # TODO(weberlo): we should probably use a more distinctive suffix for microTVM object files - elif path.endswith(".obj"): - fmt = "micro_dev" # Redirect to the load API return _ffi_api.ModuleLoadFromFile(path, fmt) diff --git a/python/tvm/runtime/ndarray.py b/python/tvm/runtime/ndarray.py index e19221c9f186..5a7acf0d6c30 100644 --- a/python/tvm/runtime/ndarray.py +++ b/python/tvm/runtime/ndarray.py @@ -268,13 +268,10 @@ def device(dev_type, dev_id=0): assert tvm.device("cuda", 0) == tvm.cuda(0) """ if isinstance(dev_type, string_types): - if "-device=micro_dev" in dev_type: - dev_type = Device.STR2MASK["micro_dev"] - else: - dev_type = dev_type.split()[0] - if dev_type not in Device.STR2MASK: - raise ValueError("Unknown device type %s" % dev_type) - dev_type = Device.STR2MASK[dev_type] + dev_type = dev_type.split()[0] + if dev_type not in Device.STR2MASK: + raise ValueError("Unknown device type %s" % dev_type) + dev_type = Device.STR2MASK[dev_type] return Device(dev_type, dev_id) @@ -510,22 +507,6 @@ def ext_dev(dev_id=0): return Device(12, dev_id) -def micro_dev(dev_id=0): - """Construct a micro device - - Parameters - ---------- - dev_id : int, optional - The integer device id - - Returns - ------- - dev : Device - The created device - """ - return Device(13, dev_id) - - def hexagon(dev_id=0): """Construct a Hexagon device diff --git a/src/runtime/module.cc b/src/runtime/module.cc index 15b9c0dde877..acc7fc7286d1 100644 --- a/src/runtime/module.cc +++ b/src/runtime/module.cc @@ -139,8 +139,6 @@ bool RuntimeEnabled(const std::string& target) { f_name = "target.build.stackvm"; } else if (target == "rpc") { f_name = "device_api.rpc"; - } else if (target == "micro_dev") { - f_name = "device_api.micro_dev"; } else if (target == "hexagon") { f_name = "device_api.hexagon"; } else if (target.length() >= 5 && target.substr(0, 5) == "nvptx") { diff --git a/src/target/codegen.cc b/src/target/codegen.cc index 19b7ad7b1d8f..cf400d90747b 100644 --- a/src/target/codegen.cc +++ b/src/target/codegen.cc @@ -47,13 +47,9 @@ runtime::Module Build(IRModule mod, Target target) { .value()) { mod = tir::transform::SkipAssert()(mod); } - std::string build_f_name; - if (target->kind->name == "micro_dev") { - build_f_name = "target.build.c"; - } else { - build_f_name = "target.build." + target->kind->name; - } + // the build function. + std::string build_f_name = "target.build." + target->kind->name; const PackedFunc* bf = runtime::Registry::Get(build_f_name); ICHECK(bf != nullptr) << build_f_name << " is not enabled"; return (*bf)(mod, target); diff --git a/src/target/source/codegen_aocl.cc b/src/target/source/codegen_aocl.cc index b3ed7cf32f7f..17e38e9af6e6 100644 --- a/src/target/source/codegen_aocl.cc +++ b/src/target/source/codegen_aocl.cc @@ -84,7 +84,7 @@ TVM_REGISTER_GLOBAL("target.build.aocl") return BuildAOCL(mod, target, false); }); -TVM_REGISTER_GLOBAL("target.build.build.aocl_sw_emu") +TVM_REGISTER_GLOBAL("target.build.aocl_sw_emu") .set_body_typed([](IRModule mod, Target target) -> runtime::Module { return BuildAOCL(mod, target, true); }); diff --git a/tests/micro/test_runtime_micro_on_arm.py b/tests/micro/test_runtime_micro_on_arm.py deleted file mode 100644 index 0212c3ea2692..000000000000 --- a/tests/micro/test_runtime_micro_on_arm.py +++ /dev/null @@ -1,370 +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. -import os - -import numpy as np -import tvm -from tvm import te -from tvm.contrib import graph_executor, utils -from tvm import relay -import tvm.micro as micro -from tvm.micro import create_micro_mod - -# Use real micro device - an STM32F746 discovery board -# SETUP: -# Be sure to have openocd installed and running -# Ex : openocd -f board/stm32f7discovery.cfg -# Be sure to have the ST CMSIS library downloaded, installed and -# Ex : export CMSIS_ST_PATH="/home/yourid/st/STM32Cube_FW_F7_V1.16.0/Drivers/CMSIS" -DEV_CONFIG_A = micro.device.arm.stm32f746xx.generate_config("127.0.0.1", 6666) -DEV_CONFIG_B = micro.device.arm.stm32f746xx.generate_config("127.0.0.1", 6666) -TARGET = "micro_dev" - - -def relay_micro_build(func, dev_config, params=None): - """Create a graph executor module with a micro device context from a Relay function. - - Parameters - ---------- - func : relay.Function - function to compile - - dev_config : Dict[str, Any] - MicroTVM config dict for the target device - - params : dict - input parameters that do not change during inference - - Return - ------ - mod : tvm.runtime.Module - graph executor module for the target device - """ - with tvm.transform.PassContext( - disabled_pass={"FuseOps"}, config={"tir.disable_vectorize": True} - ): - graph, c_mod, params = relay.build(func, target=TARGET, params=params) - micro_mod = micro.create_micro_mod(c_mod, dev_config) - ctx = tvm.micro_dev(0) - mod = graph_executor.create(graph, micro_mod, ctx) - mod.set_input(**params) - return mod - - -GDB_INIT_TEMPLATE = """ -layout asm -target remote localhost:{gdb_port} -set $pc = UTVMInit -break UTVMDone -""" - - -def reset_gdbinit(): - if "server_port" not in DEV_CONFIG_A: - return - try: - gdb_init_dir = os.environ["MICRO_GDB_INIT_DIR"] - except KeyError: - return - with open(f"{gdb_init_dir}/.gdbinit", "w") as f: - gdb_port = DEV_CONFIG_A["server_port"] - 3333 - f.write(GDB_INIT_TEMPLATE.format(gdb_port=gdb_port)) - - -def test_alloc(): - """Test tensor allocation on the device.""" - if not tvm.runtime.enabled("micro_dev"): - return - shape = (1024,) - dtype = "float32" - with micro.Session(DEV_CONFIG_A): - ctx = tvm.micro_dev(0) - np_tensor = np.random.uniform(size=shape).astype(dtype) - micro_tensor = tvm.nd.array(np_tensor, ctx) - tvm.testing.assert_allclose(np_tensor, micro_tensor.numpy()) - - -def test_add(): - """Test a module which performs addition.""" - if not tvm.runtime.enabled("micro_dev"): - return - shape = (1024,) - dtype = "float32" - - reset_gdbinit() - - # Construct TVM expression. - tvm_shape = tvm.runtime.convert(shape) - A = te.placeholder(tvm_shape, name="A", dtype=dtype) - B = te.placeholder(tvm_shape, name="B", dtype=dtype) - C = te.compute(A.shape, lambda *i: A(*i) + B(*i), name="C") - s = te.create_schedule(C.op) - - func_name = "fadd" - c_mod = tvm.build(s, [A, B, C], target="c", name=func_name) - - with micro.Session(DEV_CONFIG_A) as sess: - micro_mod = micro.create_micro_mod(c_mod, DEV_CONFIG_A) - micro_func = micro_mod[func_name] - ctx = tvm.micro_dev(0) - - a_np = np.random.uniform(size=shape).astype(dtype) - a = tvm.nd.array(a_np, ctx) - b_np = np.random.uniform(size=shape).astype(dtype) - b = tvm.nd.array(b_np, ctx) - c = tvm.nd.array(np.zeros(shape, dtype=dtype), ctx) - micro_func(a, b, c) - - # ensure inputs weren't corrupted - tvm.testing.assert_allclose(a.numpy(), a_np) - tvm.testing.assert_allclose(b.numpy(), b_np) - # ensure output is correct - tvm.testing.assert_allclose(c.numpy(), a.numpy() + b.numpy()) - - -def test_workspace_add(): - """Test a module which uses a workspace to compute an intermediate value.""" - if not tvm.runtime.enabled("micro_dev"): - return - shape = (1024,) - dtype = "float32" - - reset_gdbinit() - - # Construct TVM expression. - tvm_shape = tvm.runtime.convert(shape) - A = te.placeholder(tvm_shape, name="A", dtype=dtype) - B = te.placeholder(tvm_shape, name="B", dtype=dtype) - B = te.compute(A.shape, lambda *i: A(*i) + 1, name="B") - C = te.compute(A.shape, lambda *i: B(*i) + 1, name="C") - s = te.create_schedule(C.op) - - func_name = "fadd_two_workspace" - c_mod = tvm.build(s, [A, C], target="c", name=func_name) - - with micro.Session(DEV_CONFIG_A) as sess: - micro_mod = micro.create_micro_mod(c_mod, DEV_CONFIG_A) - micro_func = micro_mod[func_name] - ctx = tvm.micro_dev(0) - a_np = np.random.uniform(size=shape).astype(dtype) - a = tvm.nd.array(a_np, ctx) - c = tvm.nd.array(np.zeros(shape, dtype=dtype), ctx) - micro_func(a, c) - - # ensure input wasn't corrupted - tvm.testing.assert_allclose(a.numpy(), a_np) - # ensure output is correct - tvm.testing.assert_allclose(c.numpy(), a.numpy() + 2.0) - - -def test_graph_executor(): - """Test a program which uses the graph executor.""" - if not tvm.runtime.enabled("micro_dev"): - return - shape = (1024,) - dtype = "float32" - - # Construct Relay program. - x = relay.var("x", relay.TensorType(shape=shape, dtype=dtype)) - xx = relay.multiply(x, x) - z = relay.add(xx, relay.const(1.0)) - func = relay.Function([x], z) - - with micro.Session(DEV_CONFIG_A): - mod = relay_micro_build(func, DEV_CONFIG_A) - - x_in = np.random.uniform(size=shape[0]).astype(dtype) - mod.run(x=x_in) - result = mod.get_output(0).numpy() - - tvm.testing.assert_allclose(mod.get_input(0).numpy(), x_in) - tvm.testing.assert_allclose(result, x_in * x_in + 1.0) - - -def test_conv2d(): - if not tvm.runtime.enabled("micro_dev"): - return - - from tvm.relay import create_executor - from tvm.relay import transform - - dshape = (1, 4, 16, 16) - dtype = "int8" - func_name = "fused_nn_conv2d" - - reset_gdbinit() - - # Construct Relay program. - x = relay.var("x", shape=dshape, dtype=dtype) - conv_expr = relay.nn.conv2d(x, relay.var("w"), kernel_size=(3, 3), padding=(1, 1), channels=4) - func = relay.Function(relay.analysis.free_vars(conv_expr), conv_expr) - mod = tvm.IRModule.from_expr(func) - mod = transform.InferType()(mod) - - x_shape = list(map(lambda x: x.value, mod["main"].params[0].checked_type.shape)) - w_shape = list(map(lambda x: x.value, mod["main"].params[1].checked_type.shape)) - out_shape = list(map(lambda x: x.value, mod["main"].ret_type.shape)) - - with tvm.transform.PassContext(config={"tir.disable_vectorize": True}): - graph, c_mod, params = relay.build(mod, target="c") - - with micro.Session(DEV_CONFIG_A): - micro_mod = micro.create_micro_mod(c_mod, DEV_CONFIG_A) - candidate_func_name = func_name - for i in range(100): - try: - micro_func = micro_mod[candidate_func_name] - break - except tvm.TVMError as e: - candidate_func_name = f"{func_name}_{i}" - else: - assert False - ctx = tvm.micro_dev(0) - - x_data = tvm.nd.array(np.random.uniform(size=x_shape).astype(dtype), ctx) - w_data = tvm.nd.array(np.random.uniform(size=w_shape).astype(dtype), ctx) - result = tvm.nd.array(np.zeros(shape=out_shape, dtype=dtype), ctx) - micro_func(x_data, w_data, result) - - out_data = np.zeros(out_shape, dtype=dtype) - params = {"x": x_data.numpy(), "w": w_data.numpy()} - intrp = create_executor("debug") - expected_result = intrp.evaluate(mod["main"])(x_data, w_data) - - tvm.testing.assert_allclose(result.numpy(), expected_result.numpy()) - - -def test_interleave_sessions(): - """Test closing and reopening sessions.""" - if not tvm.runtime.enabled("micro_dev"): - return - shape = (1024,) - dtype = "float32" - - # Construct Relay add program. - x = relay.var("x", relay.TensorType(shape=shape, dtype=dtype)) - ret = relay.add(x, relay.const(1.0)) - add_const_func = relay.Function([x], ret) - - sess_a = micro.Session(DEV_CONFIG_A) - sess_b = micro.Session(DEV_CONFIG_B) - with sess_a: - np_tensor_a = np.random.uniform(size=shape).astype(dtype) - micro_tensor_a = tvm.nd.array(np_tensor_a, tvm.micro_dev(0)) - with sess_b: - np_tensor_b = np.random.uniform(size=shape).astype(dtype) - micro_tensor_b = tvm.nd.array(np_tensor_b, tvm.micro_dev(0)) - with sess_a: - add_const_mod = relay_micro_build(add_const_func, DEV_CONFIG_A) - add_const_mod.run(x=micro_tensor_a) - add_result = add_const_mod.get_output(0).numpy() - tvm.testing.assert_allclose(add_result, np_tensor_a + 1.0) - with sess_b: - add_const_mod = relay_micro_build(add_const_func, DEV_CONFIG_B) - add_const_mod.run(x=micro_tensor_b) - add_result = add_const_mod.get_output(0).numpy() - tvm.testing.assert_allclose(add_result, np_tensor_b + 1.0) - - -def test_nested_sessions(): - """Test entering and exiting nested session contexts.""" - if not tvm.runtime.enabled("micro_dev"): - return - shape = (1024,) - dtype = "float32" - - # Construct Relay add program. - x = relay.var("x", relay.TensorType(shape=shape, dtype=dtype)) - ret = relay.add(x, relay.const(1.0)) - add_const_func = relay.Function([x], ret) - - sess_a = micro.Session(DEV_CONFIG_A) - sess_b = micro.Session(DEV_CONFIG_B) - with sess_a: - np_tensor_a = np.random.uniform(size=shape).astype(dtype) - micro_tensor_a = tvm.nd.array(np_tensor_a, tvm.micro_dev(0)) - with sess_b: - np_tensor_b = np.random.uniform(size=shape).astype(dtype) - micro_tensor_b = tvm.nd.array(np_tensor_b, tvm.micro_dev(0)) - add_const_mod = relay_micro_build(add_const_func, DEV_CONFIG_A) - add_const_mod.run(x=micro_tensor_a) - add_result = add_const_mod.get_output(0).numpy() - tvm.testing.assert_allclose(add_result, np_tensor_a + 1.0) - - -def test_inactive_session_use(): - """Test the use of objects allocated in a session that is no longer active.""" - if not tvm.runtime.enabled("micro_dev"): - return - shape = (1024,) - dtype = "float32" - - # Construct Relay add program. - x = relay.var("x", relay.TensorType(shape=shape, dtype=dtype)) - ret = relay.add(x, relay.const(1.0)) - add_const_func = relay.Function([x], ret) - - sess_a = micro.Session(DEV_CONFIG_A) - sess_b = micro.Session(DEV_CONFIG_B) - with sess_a: - np_tensor_a = np.random.uniform(size=shape).astype(dtype) - micro_tensor_a = tvm.nd.array(np_tensor_a, tvm.micro_dev(0)) - add_const_mod = relay_micro_build(add_const_func, DEV_CONFIG_A) - - with sess_b: - # These objects belong to `sess_a`. - add_const_mod.run(x=micro_tensor_a) - add_result = add_const_mod.get_output(0).numpy() - tvm.testing.assert_allclose(add_result, np_tensor_a + 1.0) - - -# TODO add workspace alloc/free stress test - -if __name__ == "__main__": - test_alloc() - print() - print("finished alloc test") - input("[press enter to continue]") - test_add() - print() - print("finished add test") - input("[press enter to continue]") - test_workspace_add() - print() - print("finished workspace add test") - input("[press enter to continue]") - test_graph_executor() - print() - print("finished graph executor test") - input("[press enter to continue]") - test_conv2d() - print() - print("finished conv2d test") - input("[press enter to continue]") - # disable for now as these are currently broken - # test_interleave_sessions() - # print() - # print('finished interleaved sessions test') - # input('[press enter to continue]') - # test_nested_sessions() - # print() - # print('finished nested sessions test') - # input('[press enter to continue]') - test_inactive_session_use() - print() - print("finished use inactive session test") - input("[press enter to continue]")