[Target] Use TargetNode::attrs for Target serialization (#5993)

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):

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

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")

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")

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 <typename TObjectRef>

include/tvm/target/target_id.h

@@ -60,30 +60,35 @@ class TargetIdNode : public Object {
   int device_type;
   /*! \brief Default keys of the target */
   Array<String> 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<ValueTypeInfo> key;
-    std::unique_ptr<ValueTypeInfo> val;
-  };
 
   void VisitAttrs(AttrVisitor* v) {
     v->Visit("name", &name);
     v->Visit("device_type", &device_type);
     v->Visit("default_keys", &default_keys);
   }
 
-  Map<String, ObjectRef> ParseAttrsFromRawString(const std::vector<std::string>& options);
+  Map<String, ObjectRef> ParseAttrsFromRaw(const std::vector<std::string>& options) const;
+
+  Optional<String> StringifyAttrsToRaw(const Map<String, ObjectRef>& 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<ValueTypeInfo> key;
+    std::unique_ptr<ValueTypeInfo> val;
+  };
+
   uint32_t AttrRegistryIndex() const { return index_; }
   String AttrRegistryName() const { return name; }
   /*! \brief Perform schema validation */
   void ValidateSchema(const Map<String, ObjectRef>& 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<String, ValueTypeInfo> key2vtype_;
   /*! \brief A hash table that stores the default value of each attr of the target key */

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):

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
 """
 

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. """

python/tvm/target/__init__.py

@@ -26,7 +26,7 @@
 
    The device name.
 
-- **-mtriple=<target triple>** or **-target**
+- **-mtriple=<target triple>**
 
    Specify the target triple, which is useful for cross
    compilation.

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

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()

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()

src/target/llvm/codegen_blob.cc

@@ -33,7 +33,7 @@ namespace codegen {
 std::pair<std::unique_ptr<llvm::Module>, std::shared_ptr<llvm::LLVMContext>> 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<llvm::LLVMContext>();
   std::string module_name = "devc";

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;
     }
   }
 

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
  */

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();

src/target/target.cc

@@ -43,21 +43,16 @@ Target Target::CreateTarget(const std::string& name, const std::vector<std::stri
   // tag is always empty
   target->tag = "";
   // parse attrs
-  target->attrs = id->ParseAttrsFromRawString(options);
+  target->attrs = id->ParseAttrsFromRaw(options);
   String device_name = target->GetAttr<String>("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<String> keys;
     // user provided keys
-    Array<String> keys = target->GetAttr<Array<String>>("keys").value_or({});
+    if (Optional<Array<String>> user_keys = target->GetAttr<Array<String>>("keys")) {
+      keys = std::vector<String>(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::vector<std::stri
     for (const auto& key : target->id->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<std::string> 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<String> attrs_str = id->StringifyAttrsToRaw(attrs)) {
+      os << ' ' << attrs_str.value();
+    }
+    str_repr_ = os.str();
+  }
   return str_repr_;
 }