From 67a7497ad304d128f2b0d32a16718f432e33044b Mon Sep 17 00:00:00 2001 From: Junru Shao Date: Sun, 5 Jul 2020 09:52:08 -0700 Subject: [PATCH] [Target] Use TargetNode::attrs for Target serialization (#5993) --- apps/android_camera/models/prepare_model.py | 2 +- apps/android_rpc/tests/android_rpc_test.py | 2 +- apps/ios_rpc/tests/ios_rpc_mobilenet.py | 2 +- apps/ios_rpc/tests/ios_rpc_test.py | 2 +- include/tvm/target/target.h | 3 +- include/tvm/target/target_id.h | 21 +++--- python/tvm/autotvm/record.py | 4 ++ python/tvm/exec/measure_peak.py | 2 +- python/tvm/relay/qnn/op/legalizations.py | 2 +- python/tvm/target/__init__.py | 2 +- python/tvm/target/target.py | 44 ++++-------- .../tests/test_wasm32/src/build_test_lib.py | 2 +- .../tests/test_wasm32/src/build_test_lib.py | 2 +- src/target/llvm/codegen_blob.cc | 2 +- src/target/llvm/llvm_common.cc | 9 +-- src/target/llvm/llvm_common.h | 4 +- src/target/llvm/llvm_module.cc | 2 +- src/target/target.cc | 52 ++++++++++---- src/target/target_id.cc | 70 +++++++++++++++++-- tests/cpp/build_module_test.cc | 4 +- tests/python/relay/test_pass_qnn_legalize.py | 16 ++--- .../unittest/test_target_codegen_arm.py | 4 +- .../test_target_codegen_cross_llvm.py | 4 +- .../unittest/test_target_codegen_llvm.py | 4 +- tests/python/unittest/test_target_target.py | 4 +- topi/python/topi/arm_cpu/tensor_intrin.py | 2 +- topi/recipe/conv/test_conv_int8_arm.py | 2 +- topi/recipe/gemm/android_gemm_square.py | 2 +- .../python/test_topi_bitserial_conv2d_rasp.py | 2 +- tutorials/autotvm/tune_relay_arm.py | 2 +- tutorials/autotvm/tune_relay_mobile_gpu.py | 2 +- tutorials/cross_compilation_and_rpc.py | 12 ++-- tutorials/frontend/deploy_model_on_android.py | 2 +- tutorials/frontend/deploy_model_on_rasp.py | 2 +- vta/python/vta/environment.py | 4 +- web/tests/python/prepare_test_libs.py | 2 +- web/tests/python/webgpu_rpc_test.py | 2 +- web/tests/python/websock_rpc_test.py | 2 +- 38 files changed, 190 insertions(+), 113 deletions(-) diff --git a/apps/android_camera/models/prepare_model.py b/apps/android_camera/models/prepare_model.py index 703a4656c479..ab1ed78dcbed 100644 --- a/apps/android_camera/models/prepare_model.py +++ b/apps/android_camera/models/prepare_model.py @@ -28,7 +28,7 @@ from tvm.contrib import util, ndk, graph_runtime as runtime from tvm.contrib.download import download_testdata, download -target = 'llvm -target=arm64-linux-android' +target = 'llvm -mtriple=arm64-linux-android' target_host = None def del_dir(target: Union[Path, str], only_if_empty: bool = False): diff --git a/apps/android_rpc/tests/android_rpc_test.py b/apps/android_rpc/tests/android_rpc_test.py index 32af005d7d4d..754d09234fc5 100644 --- a/apps/android_rpc/tests/android_rpc_test.py +++ b/apps/android_rpc/tests/android_rpc_test.py @@ -36,7 +36,7 @@ # Change target configuration. # Run `adb shell cat /proc/cpuinfo` to find the arch. arch = "arm64" -target = "llvm -target=%s-linux-android" % arch +target = "llvm -mtriple=%s-linux-android" % arch # whether enable to execute test on OpenCL target test_opencl = False diff --git a/apps/ios_rpc/tests/ios_rpc_mobilenet.py b/apps/ios_rpc/tests/ios_rpc_mobilenet.py index e8f81ffddcec..1ce3651efdf8 100644 --- a/apps/ios_rpc/tests/ios_rpc_mobilenet.py +++ b/apps/ios_rpc/tests/ios_rpc_mobilenet.py @@ -52,7 +52,7 @@ #sdk = "iphonesimulator" arch = "arm64" sdk = "iphoneos" -target_host = "llvm -target=%s-apple-darwin" % arch +target_host = "llvm -mtriple=%s-apple-darwin" % arch # override metal compiler to compile to iphone @tvm.register_func("tvm_callback_metal_compile") diff --git a/apps/ios_rpc/tests/ios_rpc_test.py b/apps/ios_rpc/tests/ios_rpc_test.py index 973c252be175..181f3c04bbd1 100644 --- a/apps/ios_rpc/tests/ios_rpc_test.py +++ b/apps/ios_rpc/tests/ios_rpc_test.py @@ -46,7 +46,7 @@ # Change target configuration, this is setting for iphone6s arch = "arm64" sdk = "iphoneos" -target = "llvm -target=%s-apple-darwin" % arch +target = "llvm -mtriple=%s-apple-darwin" % arch # override metal compiler to compile to iphone @tvm.register_func("tvm_callback_metal_compile") diff --git a/include/tvm/target/target.h b/include/tvm/target/target.h index 30ae19ad2ab2..618095f09f6b 100644 --- a/include/tvm/target/target.h +++ b/include/tvm/target/target.h @@ -58,9 +58,8 @@ class TargetNode : public Object { void VisitAttrs(AttrVisitor* v) { v->Visit("id", &id); v->Visit("tag", &tag); - v->Visit("keys_", &keys); + v->Visit("keys", &keys); v->Visit("attrs", &attrs); - v->Visit("_str_repr_", &str_repr_); } template diff --git a/include/tvm/target/target_id.h b/include/tvm/target/target_id.h index e8d53a322eee..a0f275ce402a 100644 --- a/include/tvm/target/target_id.h +++ b/include/tvm/target/target_id.h @@ -60,13 +60,6 @@ class TargetIdNode : public Object { int device_type; /*! \brief Default keys of the target */ Array default_keys; - /*! \brief Stores the required type_key and type_index of a specific attr of a target */ - struct ValueTypeInfo { - String type_key; - uint32_t type_index; - std::unique_ptr key; - std::unique_ptr val; - }; void VisitAttrs(AttrVisitor* v) { v->Visit("name", &name); @@ -74,16 +67,28 @@ class TargetIdNode : public Object { v->Visit("default_keys", &default_keys); } - Map ParseAttrsFromRawString(const std::vector& options); + Map ParseAttrsFromRaw(const std::vector& options) const; + + Optional StringifyAttrsToRaw(const Map& attrs) const; static constexpr const char* _type_key = "TargetId"; TVM_DECLARE_FINAL_OBJECT_INFO(TargetIdNode, Object); private: + /*! \brief Stores the required type_key and type_index of a specific attr of a target */ + struct ValueTypeInfo { + String type_key; + uint32_t type_index; + std::unique_ptr key; + std::unique_ptr val; + }; + uint32_t AttrRegistryIndex() const { return index_; } String AttrRegistryName() const { return name; } /*! \brief Perform schema validation */ void ValidateSchema(const Map& config) const; + /*! \brief Verify if the obj is consistent with the type info */ + void VerifyTypeInfo(const ObjectRef& obj, const TargetIdNode::ValueTypeInfo& info) const; /*! \brief A hash table that stores the type information of each attr of the target key */ std::unordered_map key2vtype_; /*! \brief A hash table that stores the default value of each attr of the target key */ diff --git a/python/tvm/autotvm/record.py b/python/tvm/autotvm/record.py index 416b2cd57eb6..260edbb183fc 100644 --- a/python/tvm/autotvm/record.py +++ b/python/tvm/autotvm/record.py @@ -147,6 +147,10 @@ def decode(row, protocol='json'): return None tgt, task_name, task_args, task_kwargs = row["input"] + tgt = str(tgt) + if "-target" in tgt: + logger.warning("\"-target\" is deprecated, use \"-mtriple\" instead.") + tgt = tgt.replace("-target", "-mtriple") tgt = _target.create(str(tgt)) def clean_json_to_python(x): diff --git a/python/tvm/exec/measure_peak.py b/python/tvm/exec/measure_peak.py index c9759bfd2b69..2d42ed9c3e14 100644 --- a/python/tvm/exec/measure_peak.py +++ b/python/tvm/exec/measure_peak.py @@ -18,7 +18,7 @@ e.g. python3 -m tvm.exec.measure_peak --target cuda --rpc-host 0.0.0.0 --rpc-port 9090 -python3 -m tvm.exec.measure_peak --target opencl --target-host "llvm -target=aarch64-linux-gnu" \ +python3 -m tvm.exec.measure_peak --target opencl --target-host "llvm -mtriple=aarch64-linux-gnu" \ --rpc-host $TVM_OPENCL_DEVICE_HOST --rpc-port 9090 """ diff --git a/python/tvm/relay/qnn/op/legalizations.py b/python/tvm/relay/qnn/op/legalizations.py index 00866e0e3230..4d515dea329f 100644 --- a/python/tvm/relay/qnn/op/legalizations.py +++ b/python/tvm/relay/qnn/op/legalizations.py @@ -234,7 +234,7 @@ def is_fast_int8_on_intel(): def is_fast_int8_on_arm(): """ Checks whether the hardware has support for fast Int8 arithmetic operations. """ target = tvm.target.Target.current(allow_none=False) - return '+v8.2a,+dotprod' in target.mattr + return "+v8.2a" in target.mattr and "+dotprod" in target.mattr def is_aarch64_arm(): """ Checks whether we are compiling for an AArch64 target. """ diff --git a/python/tvm/target/__init__.py b/python/tvm/target/__init__.py index 18a9e7e2de7a..55bb110abf18 100644 --- a/python/tvm/target/__init__.py +++ b/python/tvm/target/__init__.py @@ -26,7 +26,7 @@ The device name. -- **-mtriple=** or **-target** +- **-mtriple=** Specify the target triple, which is useful for cross compilation. diff --git a/python/tvm/target/target.py b/python/tvm/target/target.py index a2a4501da57e..dc8c446fb6e5 100644 --- a/python/tvm/target/target.py +++ b/python/tvm/target/target.py @@ -43,19 +43,6 @@ class Target(Object): - :py:func:`tvm.target.mali` create Mali target - :py:func:`tvm.target.intel_graphics` create Intel Graphics target """ - def __new__(cls): - # Always override new to enable class - obj = Object.__new__(cls) - obj._keys = None - obj._libs = None - return obj - - @property - def keys(self): - if not self._keys: - self._keys = [str(k) for k in self.keys_] - return self._keys - def __enter__(self): _ffi_api.EnterTargetScope(self) return self @@ -103,14 +90,11 @@ def mcpu(self): @property def mattr(self): """Returns the mattr from the target if it exists.""" - return self.attrs.get("mattr", "") + return list(self.attrs.get("mattr", [])) @property def libs(self): - if not self._libs: - self._libs = list(self.attrs.get("libs", "")) - return self._libs - + return list(self.attrs.get("libs", [])) def _merge_opts(opts, new_opts): @@ -194,16 +178,16 @@ def arm_cpu(model='unknown', options=None): Additional options """ trans_table = { - "pixel2": ["-model=snapdragon835", "-target=arm64-linux-android -mattr=+neon"], - "mate10": ["-model=kirin970", "-target=arm64-linux-android -mattr=+neon"], - "mate10pro": ["-model=kirin970", "-target=arm64-linux-android -mattr=+neon"], - "p20": ["-model=kirin970", "-target=arm64-linux-android -mattr=+neon"], - "p20pro": ["-model=kirin970", "-target=arm64-linux-android -mattr=+neon"], - "rasp3b": ["-model=bcm2837", "-target=armv7l-linux-gnueabihf -mattr=+neon"], - "rasp4b": ["-model=bcm2711", "-target=arm-linux-gnueabihf -mattr=+neon"], - "rk3399": ["-model=rk3399", "-target=aarch64-linux-gnu -mattr=+neon"], - "pynq": ["-model=pynq", "-target=armv7a-linux-eabi -mattr=+neon"], - "ultra96": ["-model=ultra96", "-target=aarch64-linux-gnu -mattr=+neon"], + "pixel2": ["-model=snapdragon835", "-mtriple=arm64-linux-android -mattr=+neon"], + "mate10": ["-model=kirin970", "-mtriple=arm64-linux-android -mattr=+neon"], + "mate10pro": ["-model=kirin970", "-mtriple=arm64-linux-android -mattr=+neon"], + "p20": ["-model=kirin970", "-mtriple=arm64-linux-android -mattr=+neon"], + "p20pro": ["-model=kirin970", "-mtriple=arm64-linux-android -mattr=+neon"], + "rasp3b": ["-model=bcm2837", "-mtriple=armv7l-linux-gnueabihf -mattr=+neon"], + "rasp4b": ["-model=bcm2711", "-mtriple=arm-linux-gnueabihf -mattr=+neon"], + "rk3399": ["-model=rk3399", "-mtriple=aarch64-linux-gnu -mattr=+neon"], + "pynq": ["-model=pynq", "-mtriple=armv7a-linux-eabi -mattr=+neon"], + "ultra96": ["-model=ultra96", "-mtriple=aarch64-linux-gnu -mattr=+neon"], } pre_defined_opt = trans_table.get(model, ["-model=%s" % model]) @@ -262,7 +246,7 @@ def hexagon(cpu_ver='v66', sim_args=None, hvx=128): Size of hvx register. Value of 0 indicates disabled hvx. """ # Example compiler arguments - # llvm -target=hexagon -mcpu=hexagonv66 -mattr=+hvxv66,+hvx-length128b + # llvm -mtriple=hexagon -mcpu=hexagonv66 -mattr=+hvxv66,+hvx-length128b # Check for valid codegen cpu valid_hex = ['v60', 'v62', 'v65', 'v66', 'v67', 'v67t'] @@ -277,7 +261,7 @@ def hexagon(cpu_ver='v66', sim_args=None, hvx=128): # Target string def create_target(cpu_ver): - target = ' -target=hexagon' + target = ' -mtriple=hexagon' mcpu = ' -mcpu=hexagon' + cpu_ver mattr = '' # HVX enable diff --git a/rust/runtime/tests/test_wasm32/src/build_test_lib.py b/rust/runtime/tests/test_wasm32/src/build_test_lib.py index 6016c60c4ea3..e598bde2940c 100755 --- a/rust/runtime/tests/test_wasm32/src/build_test_lib.py +++ b/rust/runtime/tests/test_wasm32/src/build_test_lib.py @@ -32,7 +32,7 @@ def main(): s = tvm.te.create_schedule(C.op) s[C].parallel(s[C].op.axis[0]) print(tvm.lower(s, [A, B, C], simple_mode=True)) - tvm.build(s, [A, B, C], 'llvm -target=wasm32-unknown-unknown --system-lib').save(osp.join(sys.argv[1], 'test.o')) + tvm.build(s, [A, B, C], 'llvm -mtriple=wasm32-unknown-unknown --system-lib').save(osp.join(sys.argv[1], 'test.o')) if __name__ == '__main__': main() diff --git a/rust/tvm-graph-rt/tests/test_wasm32/src/build_test_lib.py b/rust/tvm-graph-rt/tests/test_wasm32/src/build_test_lib.py index 6016c60c4ea3..e598bde2940c 100755 --- a/rust/tvm-graph-rt/tests/test_wasm32/src/build_test_lib.py +++ b/rust/tvm-graph-rt/tests/test_wasm32/src/build_test_lib.py @@ -32,7 +32,7 @@ def main(): s = tvm.te.create_schedule(C.op) s[C].parallel(s[C].op.axis[0]) print(tvm.lower(s, [A, B, C], simple_mode=True)) - tvm.build(s, [A, B, C], 'llvm -target=wasm32-unknown-unknown --system-lib').save(osp.join(sys.argv[1], 'test.o')) + tvm.build(s, [A, B, C], 'llvm -mtriple=wasm32-unknown-unknown --system-lib').save(osp.join(sys.argv[1], 'test.o')) if __name__ == '__main__': main() diff --git a/src/target/llvm/codegen_blob.cc b/src/target/llvm/codegen_blob.cc index b7c48c779073..af2562ed7868 100644 --- a/src/target/llvm/codegen_blob.cc +++ b/src/target/llvm/codegen_blob.cc @@ -33,7 +33,7 @@ namespace codegen { std::pair, std::shared_ptr> CodeGenBlob( const std::string& data, bool system_lib, const std::string& target_triple) { InitializeLLVM(); - auto tm = GetLLVMTargetMachine(std::string("-target ") + target_triple); + auto tm = GetLLVMTargetMachine(std::string("-mtriple ") + target_triple); auto triple = tm->getTargetTriple(); auto ctx = std::make_shared(); std::string module_name = "devc"; diff --git a/src/target/llvm/llvm_common.cc b/src/target/llvm/llvm_common.cc index 5534a643676c..3a1036b3b0b5 100644 --- a/src/target/llvm/llvm_common.cc +++ b/src/target/llvm/llvm_common.cc @@ -73,9 +73,8 @@ void ParseLLVMTargetOptions(const std::string& target_str, std::string* triple, bool soft_float_abi = false; std::string key, value; std::istringstream is(target_str.substr(start, target_str.length() - start)); - while (is >> key) { - if (key == "--system-lib" || key == "-system-lib") { + if (key == "-system-lib" || key == "-system-lib=0" || key == "-system-lib=1") { continue; } size_t pos = key.find('='); @@ -86,7 +85,7 @@ void ParseLLVMTargetOptions(const std::string& target_str, std::string* triple, } else { CHECK(is >> value) << "Unspecified value for option " << key; } - if (key == "-target" || key == "-mtriple") { + if (key == "-mtriple") { *triple = value; } else if (key == "-mcpu") { *mcpu = value; @@ -103,10 +102,6 @@ void ParseLLVMTargetOptions(const std::string& target_str, std::string* triple, } else { LOG(FATAL) << "invalid -mfloat-abi option " << value; } - } else if (key == "-device" || key == "-libs" || key == "-model") { - // pass - } else { - LOG(FATAL) << "unknown option " << key; } } diff --git a/src/target/llvm/llvm_common.h b/src/target/llvm/llvm_common.h index 9a4ccfc9cf9c..738e0558da85 100644 --- a/src/target/llvm/llvm_common.h +++ b/src/target/llvm/llvm_common.h @@ -89,7 +89,7 @@ void InitializeLLVM(); /*! * \brief Parse target options - * \param target_str Target string, in format "llvm -target=xxx -mcpu=xxx" + * \param target_str Target string, in format "llvm -mtriple=xxx -mcpu=xxx" * \param triple Target triple * \param mcpu cpu info * \param options the options @@ -100,7 +100,7 @@ void ParseLLVMTargetOptions(const std::string& target_str, std::string* triple, /*! * \brief Get target machine from target_str string. - * \param target_str Target string, in format "llvm -target=xxx -mcpu=xxx" + * \param target_str Target string, in format "llvm -mtriple=xxx -mcpu=xxx" * \param allow_null Whether allow null to be returned. * \return target machine */ diff --git a/src/target/llvm/llvm_module.cc b/src/target/llvm/llvm_module.cc index 1151b33536b5..a3e2123568af 100644 --- a/src/target/llvm/llvm_module.cc +++ b/src/target/llvm/llvm_module.cc @@ -251,7 +251,7 @@ class LLVMModuleNode final : public runtime::ModuleNode { target_ = pstr->getString().str(); } else { std::ostringstream os; - os << "llvm -target " << module_->getTargetTriple(); + os << "llvm -mtriple " << module_->getTargetTriple(); target_ = os.str(); } mptr_ = module_.get(); diff --git a/src/target/target.cc b/src/target/target.cc index 5c6186729ff4..fb9c34e0e00c 100644 --- a/src/target/target.cc +++ b/src/target/target.cc @@ -43,21 +43,16 @@ Target Target::CreateTarget(const std::string& name, const std::vectortag = ""; // parse attrs - target->attrs = id->ParseAttrsFromRawString(options); + target->attrs = id->ParseAttrsFromRaw(options); String device_name = target->GetAttr("device", "").value(); - // create string representation - { - std::ostringstream str_repr; - str_repr << name; - for (const auto& s : options) { - str_repr << ' ' << s; - } - target->str_repr_ = str_repr.str(); - } // set up keys { + std::vector keys; // user provided keys - Array keys = target->GetAttr>("keys").value_or({}); + if (Optional> user_keys = target->GetAttr>("keys")) { + keys = std::vector(user_keys.value().begin(), user_keys.value().end()); + target->attrs.erase("keys"); + } // add `device_name` if (!device_name.empty()) { keys.push_back(device_name); @@ -66,6 +61,20 @@ Target Target::CreateTarget(const std::string& name, const std::vectorid->default_keys) { keys.push_back(key); } + // de-duplicate keys + size_t new_size = 0; + for (size_t i = 0; i < keys.size(); ++i) { + if (keys[i] == "") { + continue; + } + keys[new_size++] = keys[i]; + for (size_t j = i + 1; j < keys.size(); ++j) { + if (keys[j] == keys[i]) { + keys[j] = String(""); + } + } + } + keys.resize(new_size); target->keys = std::move(keys); } return Target(target); @@ -116,7 +125,26 @@ std::unordered_set TargetNode::GetLibs() const { } const std::string& TargetNode::str() const { - CHECK(!str_repr_.empty()); + if (str_repr_.empty()) { + std::ostringstream os; + os << id->name; + if (!this->keys.empty()) { + os << " -keys="; + bool is_first = true; + for (const String& s : keys) { + if (is_first) { + is_first = false; + } else { + os << ','; + } + os << s; + } + } + if (Optional attrs_str = id->StringifyAttrsToRaw(attrs)) { + os << ' ' << attrs_str.value(); + } + str_repr_ = os.str(); + } return str_repr_; } diff --git a/src/target/target_id.cc b/src/target/target_id.cc index faecf036015f..1841ee853d06 100644 --- a/src/target/target_id.cc +++ b/src/target/target_id.cc @@ -23,6 +23,8 @@ */ #include +#include + #include "../node/attr_registry.h" #include "../runtime/object_internal.h" @@ -57,7 +59,8 @@ const TargetId& TargetId::Get(const String& target_id_name) { return reg->id_; } -void VerifyTypeInfo(const ObjectRef& obj, const TargetIdNode::ValueTypeInfo& info) { +void TargetIdNode::VerifyTypeInfo(const ObjectRef& obj, + const TargetIdNode::ValueTypeInfo& info) const { CHECK(obj.defined()) << "Object is None"; if (!runtime::ObjectInternal::DerivedFrom(obj.get(), info.type_index)) { LOG(FATAL) << "AttributeError: expect type \"" << info.type_key << "\" but get " @@ -205,8 +208,30 @@ static inline ObjectRef ParseScalar(uint32_t type_index, const std::string& str) return ObjectRef(nullptr); } -Map TargetIdNode::ParseAttrsFromRawString( - const std::vector& options) { +static inline Optional StringifyScalar(const ObjectRef& obj) { + if (const auto* p = obj.as()) { + return String(std::to_string(p->value)); + } + if (const auto* p = obj.as()) { + return GetRef(p); + } + return NullOpt; +} + +static inline Optional Join(const std::vector& array, char separator) { + if (array.empty()) { + return NullOpt; + } + std::ostringstream os; + os << array[0]; + for (size_t i = 1; i < array.size(); ++i) { + os << separator << array[i]; + } + return String(os.str()); +} + +Map TargetIdNode::ParseAttrsFromRaw( + const std::vector& options) const { std::unordered_map attrs; for (size_t iter = 0, end = options.size(); iter < end;) { std::string s = options[iter++]; @@ -245,6 +270,9 @@ Map TargetIdNode::ParseAttrsFromRawString( } LOG(FATAL) << os.str(); } + // check if `name` has been set once + CHECK(!attrs.count(name)) << "AttributeError: key \"" << name + << "\" appears more than once in the target string"; // then `name` is valid, let's parse them // only several types are supported when parsing raw string const auto& info = it->second; @@ -285,6 +313,39 @@ Map TargetIdNode::ParseAttrsFromRawString( return attrs; } +Optional TargetIdNode::StringifyAttrsToRaw(const Map& attrs) const { + std::ostringstream os; + std::vector keys; + for (const auto& kv : attrs) { + keys.push_back(kv.first); + } + std::sort(keys.begin(), keys.end()); + std::vector result; + for (const auto& key : keys) { + const ObjectRef& obj = attrs[key]; + Optional value = NullOpt; + if (const auto* array = obj.as()) { + std::vector items; + for (const ObjectRef& item : *array) { + Optional str = StringifyScalar(item); + if (str.defined()) { + items.push_back(str.value()); + } else { + items.clear(); + break; + } + } + value = Join(items, ','); + } else { + value = StringifyScalar(obj); + } + if (value.defined()) { + result.push_back("-" + key + "=" + value.value()); + } + } + return Join(result, ' '); +} + // TODO(@junrushao1994): remove some redundant attributes TVM_REGISTER_TARGET_ID("llvm") @@ -294,9 +355,8 @@ TVM_REGISTER_TARGET_ID("llvm") .add_attr_option("model") .add_attr_option("system-lib") .add_attr_option("mcpu") - .add_attr_option("mattr") + .add_attr_option>("mattr") .add_attr_option("mtriple") - .add_attr_option("target") // FIXME: rename to mtriple .set_default_keys({"cpu"}) .set_device_type(kDLCPU); diff --git a/tests/cpp/build_module_test.cc b/tests/cpp/build_module_test.cc index fc9edf8ababe..206470f2e136 100644 --- a/tests/cpp/build_module_test.cc +++ b/tests/cpp/build_module_test.cc @@ -56,7 +56,9 @@ TEST(BuildModule, Basic) { auto module = build(lowered, target, Target()); auto mali_target = Target::Create("opencl -model=Mali-T860MP4@800Mhz -device=mali"); - CHECK_EQ(mali_target->str(), "opencl -model=Mali-T860MP4@800Mhz -device=mali"); + CHECK_EQ( + mali_target->str(), + "opencl -keys=mali,opencl,gpu -device=mali -max_num_threads=256 -model=Mali-T860MP4@800Mhz"); } TEST(BuildModule, Heterogeneous) { diff --git a/tests/python/relay/test_pass_qnn_legalize.py b/tests/python/relay/test_pass_qnn_legalize.py index 5f7deff23b06..c50c7c491c35 100644 --- a/tests/python/relay/test_pass_qnn_legalize.py +++ b/tests/python/relay/test_pass_qnn_legalize.py @@ -133,7 +133,7 @@ def _get_mod(data_dtype, kernel_dtype): assert 'cast' in legalized_mod.astext() and "qnn.conv2d" in legalized_mod.astext() # Since same dtype, there should not be any transformation - with tvm.target.create('llvm -device=arm_cpu -target=aarch64-linux-gnu -mattr=+v8.2a,+dotprod'): + with tvm.target.create('llvm -device=arm_cpu -mtriple=aarch64-linux-gnu -mattr=+v8.2a,+dotprod'): legalized_mod = relay.qnn.transform.Legalize()(mod) assert tvm.ir.structural_equal(mod, legalized_mod) @@ -146,7 +146,7 @@ def _get_mod(data_dtype, kernel_dtype): assert 'cast' in legalized_mod.astext() and "qnn" not in legalized_mod.astext() # Older ARM vesions. - with tvm.target.create('llvm -device=arm_cpu -target=aarch64-linux-gnu'): + with tvm.target.create('llvm -device=arm_cpu -mtriple=aarch64-linux-gnu'): legalized_mod = relay.qnn.transform.Legalize()(mod) assert 'cast' in legalized_mod.astext() and "qnn" not in legalized_mod.astext() @@ -161,7 +161,7 @@ def _get_mod(data_dtype, kernel_dtype): assert tvm.ir.structural_equal(mod, legalized_mod) # ARM - so check that transformation has happened. - with tvm.target.create('llvm -device=arm_cpu -target=aarch64-linux-gnu -mattr=+v8.2a,+dotprod'): + with tvm.target.create('llvm -device=arm_cpu -mtriple=aarch64-linux-gnu -mattr=+v8.2a,+dotprod'): legalized_mod = relay.qnn.transform.Legalize()(mod) assert 'cast' in legalized_mod.astext() and "qnn.conv2d" in legalized_mod.astext() @@ -174,7 +174,7 @@ def _get_mod(data_dtype, kernel_dtype): assert 'cast' in legalized_mod.astext() and "qnn" not in legalized_mod.astext() # Older ARM vesions. - with tvm.target.create('llvm -device=arm_cpu -target=aarch64-linux-gnu'): + with tvm.target.create('llvm -device=arm_cpu -mtriple=aarch64-linux-gnu'): legalized_mod = relay.qnn.transform.Legalize()(mod) assert 'cast' in legalized_mod.astext() and "qnn" not in legalized_mod.astext() @@ -220,7 +220,7 @@ def _get_mod(data_dtype, kernel_dtype): assert 'cast' in legalized_mod.astext() and "qnn.dense" in legalized_mod.astext() # Since same dtype, there should not be any transformation - with tvm.target.create('llvm -device=arm_cpu -target=aarch64-linux-gnu -mattr=+v8.2a,+dotprod'): + with tvm.target.create('llvm -device=arm_cpu -mtriple=aarch64-linux-gnu -mattr=+v8.2a,+dotprod'): legalized_mod = relay.qnn.transform.Legalize()(mod) assert tvm.ir.structural_equal(mod, legalized_mod) @@ -233,7 +233,7 @@ def _get_mod(data_dtype, kernel_dtype): assert 'cast' in legalized_mod.astext() and "qnn" not in legalized_mod.astext() # Older ARM vesions. - with tvm.target.create('llvm -device=arm_cpu -target=aarch64-linux-gnu'): + with tvm.target.create('llvm -device=arm_cpu -mtriple=aarch64-linux-gnu'): legalized_mod = relay.qnn.transform.Legalize()(mod) assert 'cast' in legalized_mod.astext() and "qnn" not in legalized_mod.astext() @@ -248,7 +248,7 @@ def _get_mod(data_dtype, kernel_dtype): assert tvm.ir.structural_equal(mod, legalized_mod) # ARM - so check that transformation has happened. - with tvm.target.create('llvm -device=arm_cpu -target=aarch64-linux-gnu -mattr=+v8.2a,+dotprod'): + with tvm.target.create('llvm -device=arm_cpu -mtriple=aarch64-linux-gnu -mattr=+v8.2a,+dotprod'): legalized_mod = relay.qnn.transform.Legalize()(mod) assert 'cast' in legalized_mod.astext() and "qnn.dense" in legalized_mod.astext() @@ -261,7 +261,7 @@ def _get_mod(data_dtype, kernel_dtype): assert 'cast' in legalized_mod.astext() and "qnn" not in legalized_mod.astext() # Older ARM vesions. - with tvm.target.create('llvm -device=arm_cpu -target=aarch64-linux-gnu'): + with tvm.target.create('llvm -device=arm_cpu -mtriple=aarch64-linux-gnu'): legalized_mod = relay.qnn.transform.Legalize()(mod) assert 'cast' in legalized_mod.astext() and "qnn" not in legalized_mod.astext() diff --git a/tests/python/unittest/test_target_codegen_arm.py b/tests/python/unittest/test_target_codegen_arm.py index 65d82b0146fb..ace41d040353 100644 --- a/tests/python/unittest/test_target_codegen_arm.py +++ b/tests/python/unittest/test_target_codegen_arm.py @@ -21,7 +21,7 @@ import ctypes def test_popcount(): - target = 'llvm -target=armv7l-none-linux-gnueabihf -mcpu=cortex-a53 -mattr=+neon' + target = 'llvm -mtriple=armv7l-none-linux-gnueabihf -mcpu=cortex-a53 -mattr=+neon' def check_correct_assembly(type, elements, counts): n = tvm.runtime.convert(elements) @@ -45,7 +45,7 @@ def check_correct_assembly(type, elements, counts): def test_vmlal_s16(): - target = 'llvm -target=armv7l-none-linux-gnueabihf -mcpu=cortex-a53 -mattr=+neon' + target = 'llvm -mtriple=armv7l-none-linux-gnueabihf -mcpu=cortex-a53 -mattr=+neon' def check_correct_assembly(N): K = te.size_var("K") diff --git a/tests/python/unittest/test_target_codegen_cross_llvm.py b/tests/python/unittest/test_target_codegen_cross_llvm.py index cb3986eaf20f..3ea413c51d84 100644 --- a/tests/python/unittest/test_target_codegen_cross_llvm.py +++ b/tests/python/unittest/test_target_codegen_cross_llvm.py @@ -47,14 +47,14 @@ def build_i386(): print("Skip because llvm is not enabled..") return temp = util.tempdir() - target = "llvm -target=i386-pc-linux-gnu" + target = "llvm -mtriple=i386-pc-linux-gnu" f = tvm.build(s, [A, B, C], target) path = temp.relpath("myadd.o") f.save(path) verify_elf(path, 0x03) def build_arm(): - target = "llvm -target=armv7-none-linux-gnueabihf" + target = "llvm -mtriple=armv7-none-linux-gnueabihf" if not tvm.runtime.enabled(target): print("Skip because %s is not enabled.." % target) return diff --git a/tests/python/unittest/test_target_codegen_llvm.py b/tests/python/unittest/test_target_codegen_llvm.py index 911ffb44f353..cf8250401246 100644 --- a/tests/python/unittest/test_target_codegen_llvm.py +++ b/tests/python/unittest/test_target_codegen_llvm.py @@ -685,7 +685,7 @@ def check_llvm_ir(): # build two functions f2 = tvm.lower(s, [A, B, C], name="fadd1") f1 = tvm.lower(s, [A, B, C], name="fadd2") - m = tvm.build([f1, f2], target="llvm -target=aarch64-linux-gnu") + m = tvm.build([f1, f2], target="llvm -mtriple=aarch64-linux-gnu") ll = m.get_source("ll") # On non-Darwin OS, don't explicitly specify DWARF version. @@ -695,7 +695,7 @@ def check_llvm_ir(): # Try Darwin, require DWARF-2 m = tvm.build([f1, f2], - target="llvm -target=x86_64-apple-darwin-macho") + target="llvm -mtriple=x86_64-apple-darwin-macho") ll = m.get_source("ll") assert re.search(r"""i32 4, !"Dwarf Version", i32 2""", ll) assert re.search(r"""llvm.dbg.value""", ll) diff --git a/tests/python/unittest/test_target_target.py b/tests/python/unittest/test_target_target.py index fe3799bcb963..0fd2506bd3ba 100644 --- a/tests/python/unittest/test_target_target.py +++ b/tests/python/unittest/test_target_target.py @@ -59,8 +59,8 @@ def test_target_string_parse(): assert target.id.name == "cuda" assert target.model == "unknown" - assert target.keys == ['cuda', 'gpu'] - assert target.libs == ['cublas', 'cudnn'] + assert set(target.keys) == set(['cuda', 'gpu']) + assert set(target.libs) == set(['cublas', 'cudnn']) assert str(target) == str(tvm.target.cuda(options="-libs=cublas,cudnn")) assert tvm.target.intel_graphics().device_name == "intel_graphics" diff --git a/topi/python/topi/arm_cpu/tensor_intrin.py b/topi/python/topi/arm_cpu/tensor_intrin.py index dfa2f05e7960..d8d9481c2a32 100644 --- a/topi/python/topi/arm_cpu/tensor_intrin.py +++ b/topi/python/topi/arm_cpu/tensor_intrin.py @@ -267,7 +267,7 @@ def gemv_quantized_impl(M, N, data_type='uint8'): ll_path = temp.relpath("temp.ll") # Create LLVM ir from c source code ll_code = clang.create_llvm(cc_code, - options=["--target=aarch64-linux-gnu -mattr=+neon"], + options=["-mtriple=aarch64-linux-gnu -mattr=+neon"], output=ll_path) return ll_code diff --git a/topi/recipe/conv/test_conv_int8_arm.py b/topi/recipe/conv/test_conv_int8_arm.py index 336e2f2f405b..f0b260e18a98 100644 --- a/topi/recipe/conv/test_conv_int8_arm.py +++ b/topi/recipe/conv/test_conv_int8_arm.py @@ -58,7 +58,7 @@ ] -TARGET_NAME = 'llvm -device=arm_cpu -target=aarch64-linux-gnu -mattr=+v8.2a,+dotprod' +TARGET_NAME = 'llvm -device=arm_cpu -mtriple=aarch64-linux-gnu -mattr=+v8.2a,+dotprod' NUM_VEC_LANES = 16 CTX = tvm.context(TARGET_NAME, 0) diff --git a/topi/recipe/gemm/android_gemm_square.py b/topi/recipe/gemm/android_gemm_square.py index 7692f9cf4497..1443aea6c173 100644 --- a/topi/recipe/gemm/android_gemm_square.py +++ b/topi/recipe/gemm/android_gemm_square.py @@ -30,7 +30,7 @@ # Change target configuration. # Run `adb shell cat /proc/cpuinfo` to find the arch. arch = "arm64" -target = "llvm -target=%s-linux-android" % arch +target = "llvm -mtriple=%s-linux-android" % arch def ngflops(N): return 2.0 * float(N * N * N) / (10**9) diff --git a/topi/tests/python/test_topi_bitserial_conv2d_rasp.py b/topi/tests/python/test_topi_bitserial_conv2d_rasp.py index 99ba0dba8328..643a7d5aa814 100644 --- a/topi/tests/python/test_topi_bitserial_conv2d_rasp.py +++ b/topi/tests/python/test_topi_bitserial_conv2d_rasp.py @@ -36,7 +36,7 @@ def verify_bitserial_conv2d_nhwc(batch, in_size, in_channel, num_filter, kernel, input_type = 'uint32' out_dtype = 'int16' - device = 'llvm -device=arm_cpu -model=bcm2837 -target=armv7l-linux-gnueabihf -mattr=+neon' + device = 'llvm -device=arm_cpu -model=bcm2837 -mtriple=armv7l-linux-gnueabihf -mattr=+neon' with tvm.target.create(device): A = te.placeholder((batch, in_height, in_width, in_channel), dtype=input_type, name='A') W = te.placeholder((kernel, kernel, in_channel, num_filter), dtype=input_type, name='W') diff --git a/tutorials/autotvm/tune_relay_arm.py b/tutorials/autotvm/tune_relay_arm.py index 3b07097ce696..d7529b2d109a 100644 --- a/tutorials/autotvm/tune_relay_arm.py +++ b/tutorials/autotvm/tune_relay_arm.py @@ -190,7 +190,7 @@ def get_network(name, batch_size): # Replace "aarch64-linux-gnu" with the correct target of your board. # This target is used for cross compilation. You can query it by :code:`gcc -v` on your device. -target = tvm.target.create('llvm -device=arm_cpu -target=aarch64-linux-gnu') +target = tvm.target.create('llvm -device=arm_cpu -mtriple=aarch64-linux-gnu') # Also replace this with the device key in your tracker device_key = 'rk3399' diff --git a/tutorials/autotvm/tune_relay_mobile_gpu.py b/tutorials/autotvm/tune_relay_mobile_gpu.py index 4748f41e96c3..5f5e523cfe7a 100644 --- a/tutorials/autotvm/tune_relay_mobile_gpu.py +++ b/tutorials/autotvm/tune_relay_mobile_gpu.py @@ -191,7 +191,7 @@ def get_network(name, batch_size): # Replace "aarch64-linux-gnu" with the correct target of your board. # This target host is used for cross compilation. You can query it by :code:`gcc -v` on your device. -target_host = 'llvm -target=aarch64-linux-gnu' +target_host = 'llvm -mtriple=aarch64-linux-gnu' # Also replace this with the device key in your tracker device_key = 'rk3399' diff --git a/tutorials/cross_compilation_and_rpc.py b/tutorials/cross_compilation_and_rpc.py index 553d77dd2023..eaf6f03bee11 100644 --- a/tutorials/cross_compilation_and_rpc.py +++ b/tutorials/cross_compilation_and_rpc.py @@ -107,7 +107,7 @@ ###################################################################### # Then we cross compile the kernel. -# The target should be 'llvm -target=armv7l-linux-gnueabihf' for +# The target should be 'llvm -mtriple=armv7l-linux-gnueabihf' for # Raspberry Pi 3B, but we use 'llvm' here to make this tutorial runnable # on our webpage building server. See the detailed note in the following block. @@ -116,7 +116,7 @@ if local_demo: target = 'llvm' else: - target = 'llvm -target=armv7l-linux-gnueabihf' + target = 'llvm -mtriple=armv7l-linux-gnueabihf' func = tvm.build(s, [A, B], target=target, name='add_one') # save the lib at a local temp folder @@ -131,14 +131,14 @@ # to False and replace :code:`target` in :code:`build` with the appropriate # target triple for your device. The target triple which might be # different for different devices. For example, it is -# :code:`'llvm -target=armv7l-linux-gnueabihf'` for Raspberry Pi 3B and -# :code:`'llvm -target=aarch64-linux-gnu'` for RK3399. +# :code:`'llvm -mtriple=armv7l-linux-gnueabihf'` for Raspberry Pi 3B and +# :code:`'llvm -mtriple=aarch64-linux-gnu'` for RK3399. # # Usually, you can query the target by running :code:`gcc -v` on your # device, and looking for the line starting with :code:`Target:` # (Though it may still be a loose configuration.) # -# Besides :code:`-target`, you can also set other compilation options +# Besides :code:`-mtriple`, you can also set other compilation options # like: # # * -mcpu= @@ -224,7 +224,7 @@ def run_opencl(): # NOTE: This is the setting for my rk3399 board. You need to modify # them according to your environment. - target_host = "llvm -target=aarch64-linux-gnu" + target_host = "llvm -mtriple=aarch64-linux-gnu" opencl_device_host = '10.77.1.145' opencl_device_port = 9090 diff --git a/tutorials/frontend/deploy_model_on_android.py b/tutorials/frontend/deploy_model_on_android.py index bc5b5239a889..51ffdd5636d4 100644 --- a/tutorials/frontend/deploy_model_on_android.py +++ b/tutorials/frontend/deploy_model_on_android.py @@ -246,7 +246,7 @@ def transform_image(image): # Change target configuration. # Run `adb shell cat /proc/cpuinfo` to find the arch. arch = 'arm64' -target = 'llvm -target=%s-linux-android' % arch +target = 'llvm -mtriple=%s-linux-android' % arch target_host = None if local_demo: diff --git a/tutorials/frontend/deploy_model_on_rasp.py b/tutorials/frontend/deploy_model_on_rasp.py index 25df34128415..a97d411287ea 100644 --- a/tutorials/frontend/deploy_model_on_rasp.py +++ b/tutorials/frontend/deploy_model_on_rasp.py @@ -177,7 +177,7 @@ def transform_image(image): else: target = tvm.target.arm_cpu('rasp3b') # The above line is a simple form of - # target = tvm.target.create('llvm -device=arm_cpu -model=bcm2837 -target=armv7l-linux-gnueabihf -mattr=+neon') + # target = tvm.target.create('llvm -device=arm_cpu -model=bcm2837 -mtriple=armv7l-linux-gnueabihf -mattr=+neon') with tvm.transform.PassContext(opt_level=3): graph, lib, params = relay.build(func, target, params=params) diff --git a/vta/python/vta/environment.py b/vta/python/vta/environment.py index 3e6a0c538b12..3d7423848501 100644 --- a/vta/python/vta/environment.py +++ b/vta/python/vta/environment.py @@ -237,9 +237,9 @@ def target(self): def target_host(self): """The target host""" if self.TARGET in ["pynq", "de10nano"]: - return "llvm -target=armv7-none-linux-gnueabihf" + return "llvm -mtriple=armv7-none-linux-gnueabihf" if self.TARGET == "ultra96": - return "llvm -target=aarch64-linux-gnu" + return "llvm -mtriple=aarch64-linux-gnu" if self.TARGET in ["sim", "tsim"]: return "llvm" raise ValueError("Unknown target %s" % self.TARGET) diff --git a/web/tests/python/prepare_test_libs.py b/web/tests/python/prepare_test_libs.py index ec4eb5be1536..f529e04971f7 100644 --- a/web/tests/python/prepare_test_libs.py +++ b/web/tests/python/prepare_test_libs.py @@ -23,7 +23,7 @@ def prepare_test_libs(base_path): - target = "llvm -target=wasm32-unknown-unknown-wasm -system-lib" + target = "llvm -mtriple=wasm32-unknown-unknown-wasm -system-lib" if not tvm.runtime.enabled(target): raise RuntimeError("Target %s is not enbaled" % target) n = te.var("n") diff --git a/web/tests/python/webgpu_rpc_test.py b/web/tests/python/webgpu_rpc_test.py index d16ba3f3304e..6bef33a342e4 100644 --- a/web/tests/python/webgpu_rpc_test.py +++ b/web/tests/python/webgpu_rpc_test.py @@ -35,7 +35,7 @@ def test_rpc(): return # generate the wasm library target_device = "webgpu" - target_host = "llvm -target=wasm32-unknown-unknown-wasm -system-lib" + target_host = "llvm -mtriple=wasm32-unknown-unknown-wasm -system-lib" if not tvm.runtime.enabled(target_host): raise RuntimeError("Target %s is not enbaled" % target_host) diff --git a/web/tests/python/websock_rpc_test.py b/web/tests/python/websock_rpc_test.py index f7c07924a210..80f39a3e21c2 100644 --- a/web/tests/python/websock_rpc_test.py +++ b/web/tests/python/websock_rpc_test.py @@ -33,7 +33,7 @@ def test_rpc(): if not tvm.runtime.enabled("rpc"): return # generate the wasm library - target = "llvm -target=wasm32-unknown-unknown-wasm -system-lib" + target = "llvm -mtriple=wasm32-unknown-unknown-wasm -system-lib" if not tvm.runtime.enabled(target): raise RuntimeError("Target %s is not enbaled" % target) n = te.var("n")