#include <tvm/ir/transform.h>
#include <tvm/node/container.h>
#include <tvm/support/with.h>
-#include <tvm/target/target_id.h>
+#include <tvm/target/target_kind.h>
#include <string>
#include <unordered_set>
*/
class TargetNode : public Object {
public:
- /*! \brief The id of the target device */
- TargetId id;
+ /*! \brief The kind of the target device */
+ TargetKind kind;
/*! \brief Tag of the the target, can be empty */
String tag;
/*! \brief Keys for this target */
TVM_DLL const std::string& str() const;
void VisitAttrs(AttrVisitor* v) {
- v->Visit("id", &id);
+ v->Visit("kind", &kind);
v->Visit("tag", &tag);
v->Visit("keys", &keys);
v->Visit("attrs", &attrs);
*/
/*!
- * \file tvm/target/target_id.h
- * \brief Target id registry
+ * \file tvm/target/target_kind.h
+ * \brief Target kind registry
*/
-#ifndef TVM_TARGET_TARGET_ID_H_
-#define TVM_TARGET_TARGET_ID_H_
+#ifndef TVM_TARGET_TARGET_KIND_H_
+#define TVM_TARGET_TARGET_KIND_H_
#include <tvm/ir/expr.h>
#include <tvm/ir/transform.h>
TVM_DLL void TargetValidateSchema(const Map<String, ObjectRef>& config);
template <typename>
-class TargetIdAttrMap;
+class TargetKindAttrMap;
-/*! \brief Target Id, specifies the kind of the target */
-class TargetIdNode : public Object {
+/*! \brief Target kind, specifies the kind of the target */
+class TargetKindNode : public Object {
public:
- /*! \brief Name of the target id */
+ /*! \brief Name of the target kind */
String name;
- /*! \brief Device type of target id */
+ /*! \brief Device type of target kind */
int device_type;
/*! \brief Default keys of the target */
Array<String> default_keys;
Optional<String> StringifyAttrsToRaw(const Map<String, ObjectRef>& attrs) const;
- static constexpr const char* _type_key = "TargetId";
- TVM_DECLARE_FINAL_OBJECT_INFO(TargetIdNode, Object);
+ static constexpr const char* _type_key = "TargetKind";
+ TVM_DECLARE_FINAL_OBJECT_INFO(TargetKindNode, Object);
private:
/*! \brief Stores the required type_key and type_index of a specific attr of a target */
/*! \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;
+ void VerifyTypeInfo(const ObjectRef& obj, const TargetKindNode::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 */
uint32_t index_;
friend void TargetValidateSchema(const Map<String, ObjectRef>&);
friend class Target;
- friend class TargetId;
+ friend class TargetKind;
template <typename, typename>
friend class AttrRegistry;
template <typename>
friend class AttrRegistryMapContainerMap;
- friend class TargetIdRegEntry;
+ friend class TargetKindRegEntry;
template <typename, typename, typename>
friend struct detail::ValueTypeInfoMaker;
};
/*!
- * \brief Managed reference class to TargetIdNode
- * \sa TargetIdNode
+ * \brief Managed reference class to TargetKindNode
+ * \sa TargetKindNode
*/
-class TargetId : public ObjectRef {
+class TargetKind : public ObjectRef {
public:
- TargetId() = default;
+ TargetKind() = default;
/*! \brief Get the attribute map given the attribute name */
template <typename ValueType>
- static inline TargetIdAttrMap<ValueType> GetAttrMap(const String& attr_name);
+ static inline TargetKindAttrMap<ValueType> GetAttrMap(const String& attr_name);
/*!
- * \brief Retrieve the TargetId given its name
- * \param target_id_name Name of the target id
- * \return The TargetId requested
+ * \brief Retrieve the TargetKind given its name
+ * \param target_kind_name Name of the target kind
+ * \return The TargetKind requested
*/
- TVM_DLL static const TargetId& Get(const String& target_id_name);
- TVM_DEFINE_NOTNULLABLE_OBJECT_REF_METHODS(TargetId, ObjectRef, TargetIdNode);
+ TVM_DLL static const TargetKind& Get(const String& target_kind_name);
+ TVM_DEFINE_NOTNULLABLE_OBJECT_REF_METHODS(TargetKind, ObjectRef, TargetKindNode);
private:
/*! \brief Mutable access to the container class */
- TargetIdNode* operator->() { return static_cast<TargetIdNode*>(data_.get()); }
- TVM_DLL static const AttrRegistryMapContainerMap<TargetId>& GetAttrMapContainer(
+ TargetKindNode* operator->() { return static_cast<TargetKindNode*>(data_.get()); }
+ TVM_DLL static const AttrRegistryMapContainerMap<TargetKind>& GetAttrMapContainer(
const String& attr_name);
template <typename, typename>
friend class AttrRegistry;
- friend class TargetIdRegEntry;
+ friend class TargetKindRegEntry;
friend class Target;
};
/*!
- * \brief Map<TargetId, ValueType> used to store meta-information about TargetId
+ * \brief Map<TargetKind, ValueType> used to store meta-information about TargetKind
* \tparam ValueType The type of the value stored in map
*/
template <typename ValueType>
-class TargetIdAttrMap : public AttrRegistryMap<TargetId, ValueType> {
+class TargetKindAttrMap : public AttrRegistryMap<TargetKind, ValueType> {
public:
- using TParent = AttrRegistryMap<TargetId, ValueType>;
+ using TParent = AttrRegistryMap<TargetKind, ValueType>;
using TParent::count;
using TParent::get;
using TParent::operator[];
- explicit TargetIdAttrMap(const AttrRegistryMapContainerMap<TargetId>& map) : TParent(map) {}
+ explicit TargetKindAttrMap(const AttrRegistryMapContainerMap<TargetKind>& map) : TParent(map) {}
};
/*!
- * \brief Helper structure to register TargetId
- * \sa TVM_REGISTER_TARGET_ID
+ * \brief Helper structure to register TargetKind
+ * \sa TVM_REGISTER_TARGET_KIND
*/
-class TargetIdRegEntry {
+class TargetKindRegEntry {
public:
/*!
- * \brief Register additional attributes to target_id.
+ * \brief Register additional attributes to target_kind.
* \param attr_name The name of the attribute.
* \param value The value to be set.
* \param plevel The priority level of this attribute,
* \tparam ValueType The type of the value to be set.
*/
template <typename ValueType>
- inline TargetIdRegEntry& set_attr(const String& attr_name, const ValueType& value,
- int plevel = 10);
+ inline TargetKindRegEntry& set_attr(const String& attr_name, const ValueType& value,
+ int plevel = 10);
/*!
* \brief Set DLPack's device_type the target
* \param device_type Device type
*/
- inline TargetIdRegEntry& set_device_type(int device_type);
+ inline TargetKindRegEntry& set_device_type(int device_type);
/*!
* \brief Set DLPack's device_type the target
* \param keys The default keys
*/
- inline TargetIdRegEntry& set_default_keys(std::vector<String> keys);
+ inline TargetKindRegEntry& set_default_keys(std::vector<String> keys);
/*!
* \brief Register a valid configuration option and its ValueType for validation
* \param key The configuration key
* \tparam ValueType The value type to be registered
*/
template <typename ValueType>
- inline TargetIdRegEntry& add_attr_option(const String& key);
+ inline TargetKindRegEntry& add_attr_option(const String& key);
/*!
* \brief Register a valid configuration option and its ValueType for validation
* \param key The configuration key
* \tparam ValueType The value type to be registered
*/
template <typename ValueType>
- inline TargetIdRegEntry& add_attr_option(const String& key, ObjectRef default_value);
- /*! \brief Set name of the TargetId to be the same as registry if it is empty */
- inline TargetIdRegEntry& set_name();
+ inline TargetKindRegEntry& add_attr_option(const String& key, ObjectRef default_value);
+ /*! \brief Set name of the TargetKind to be the same as registry if it is empty */
+ inline TargetKindRegEntry& set_name();
/*!
* \brief Register or get a new entry.
- * \param target_id_name The name of the TargetId.
+ * \param target_kind_name The name of the TargetKind.
* \return the corresponding entry.
*/
- TVM_DLL static TargetIdRegEntry& RegisterOrGet(const String& target_id_name);
+ TVM_DLL static TargetKindRegEntry& RegisterOrGet(const String& target_kind_name);
private:
- TargetId id_;
+ TargetKind kind_;
String name;
/*! \brief private constructor */
- explicit TargetIdRegEntry(uint32_t reg_index) : id_(make_object<TargetIdNode>()) {
- id_->index_ = reg_index;
+ explicit TargetKindRegEntry(uint32_t reg_index) : kind_(make_object<TargetKindNode>()) {
+ kind_->index_ = reg_index;
}
/*!
- * \brief update the attribute TargetIdAttrMap
+ * \brief update the attribute TargetKindAttrMap
* \param key The name of the attribute
* \param value The value to be set
* \param plevel The priority level
TVM_DLL void UpdateAttr(const String& key, TVMRetValue value, int plevel);
template <typename, typename>
friend class AttrRegistry;
- friend class TargetId;
+ friend class TargetKind;
};
-#define TVM_TARGET_ID_REGISTER_VAR_DEF \
- static DMLC_ATTRIBUTE_UNUSED ::tvm::TargetIdRegEntry& __make_##TargetId
+#define TVM_TARGET_KIND_REGISTER_VAR_DEF \
+ static DMLC_ATTRIBUTE_UNUSED ::tvm::TargetKindRegEntry& __make_##TargetKind
/*!
- * \def TVM_REGISTER_TARGET_ID
- * \brief Register a new target id, or set attribute of the corresponding target id.
+ * \def TVM_REGISTER_TARGET_KIND
+ * \brief Register a new target kind, or set attribute of the corresponding target kind.
*
- * \param TargetIdName The name of target id
+ * \param TargetKindName The name of target kind
*
* \code
*
- * TVM_REGISTER_TARGET_ID("llvm")
+ * TVM_REGISTER_TARGET_KIND("llvm")
* .set_attr<TPreCodegenPass>("TPreCodegenPass", a-pre-codegen-pass)
* .add_attr_option<Bool>("system_lib")
* .add_attr_option<String>("mtriple")
*
* \endcode
*/
-#define TVM_REGISTER_TARGET_ID(TargetIdName) \
- TVM_STR_CONCAT(TVM_TARGET_ID_REGISTER_VAR_DEF, __COUNTER__) = \
- ::tvm::TargetIdRegEntry::RegisterOrGet(TargetIdName).set_name()
+#define TVM_REGISTER_TARGET_KIND(TargetKindName) \
+ TVM_STR_CONCAT(TVM_TARGET_KIND_REGISTER_VAR_DEF, __COUNTER__) = \
+ ::tvm::TargetKindRegEntry::RegisterOrGet(TargetKindName).set_name()
namespace detail {
template <typename Type, template <typename...> class Container>
template <typename ValueType>
struct ValueTypeInfoMaker<ValueType, std::false_type, std::false_type> {
- using ValueTypeInfo = TargetIdNode::ValueTypeInfo;
+ using ValueTypeInfo = TargetKindNode::ValueTypeInfo;
ValueTypeInfo operator()() const {
uint32_t tindex = ValueType::ContainerType::_GetOrAllocRuntimeTypeIndex();
template <typename ValueType>
struct ValueTypeInfoMaker<ValueType, std::true_type, std::false_type> {
- using ValueTypeInfo = TargetIdNode::ValueTypeInfo;
+ using ValueTypeInfo = TargetKindNode::ValueTypeInfo;
ValueTypeInfo operator()() const {
using key_type = ValueTypeInfoMaker<typename ValueType::value_type>;
template <typename ValueType>
struct ValueTypeInfoMaker<ValueType, std::false_type, std::true_type> {
- using ValueTypeInfo = TargetIdNode::ValueTypeInfo;
+ using ValueTypeInfo = TargetKindNode::ValueTypeInfo;
ValueTypeInfo operator()() const {
using key_type = ValueTypeInfoMaker<typename ValueType::key_type>;
using val_type = ValueTypeInfoMaker<typename ValueType::mapped_type>;
} // namespace detail
template <typename ValueType>
-inline TargetIdAttrMap<ValueType> TargetId::GetAttrMap(const String& attr_name) {
- return TargetIdAttrMap<ValueType>(GetAttrMapContainer(attr_name));
+inline TargetKindAttrMap<ValueType> TargetKind::GetAttrMap(const String& attr_name) {
+ return TargetKindAttrMap<ValueType>(GetAttrMapContainer(attr_name));
}
template <typename ValueType>
-inline TargetIdRegEntry& TargetIdRegEntry::set_attr(const String& attr_name, const ValueType& value,
- int plevel) {
+inline TargetKindRegEntry& TargetKindRegEntry::set_attr(const String& attr_name,
+ const ValueType& value, int plevel) {
CHECK_GT(plevel, 0) << "plevel in set_attr must be greater than 0";
runtime::TVMRetValue rv;
rv = value;
return *this;
}
-inline TargetIdRegEntry& TargetIdRegEntry::set_device_type(int device_type) {
- id_->device_type = device_type;
+inline TargetKindRegEntry& TargetKindRegEntry::set_device_type(int device_type) {
+ kind_->device_type = device_type;
return *this;
}
-inline TargetIdRegEntry& TargetIdRegEntry::set_default_keys(std::vector<String> keys) {
- id_->default_keys = keys;
+inline TargetKindRegEntry& TargetKindRegEntry::set_default_keys(std::vector<String> keys) {
+ kind_->default_keys = keys;
return *this;
}
template <typename ValueType>
-inline TargetIdRegEntry& TargetIdRegEntry::add_attr_option(const String& key) {
- CHECK(!id_->key2vtype_.count(key))
+inline TargetKindRegEntry& TargetKindRegEntry::add_attr_option(const String& key) {
+ CHECK(!kind_->key2vtype_.count(key))
<< "AttributeError: add_attr_option failed because '" << key << "' has been set once";
- id_->key2vtype_[key] = detail::ValueTypeInfoMaker<ValueType>()();
+ kind_->key2vtype_[key] = detail::ValueTypeInfoMaker<ValueType>()();
return *this;
}
template <typename ValueType>
-inline TargetIdRegEntry& TargetIdRegEntry::add_attr_option(const String& key,
- ObjectRef default_value) {
+inline TargetKindRegEntry& TargetKindRegEntry::add_attr_option(const String& key,
+ ObjectRef default_value) {
add_attr_option<ValueType>(key);
- id_->key2default_[key] = default_value;
+ kind_->key2default_[key] = default_value;
return *this;
}
-inline TargetIdRegEntry& TargetIdRegEntry::set_name() {
- if (id_->name.empty()) {
- id_->name = name;
+inline TargetKindRegEntry& TargetKindRegEntry::set_name() {
+ if (kind_->name.empty()) {
+ kind_->name = name;
}
return *this;
}
} // namespace tvm
-#endif // TVM_TARGET_TARGET_ID_H_
+#endif // TVM_TARGET_TARGET_KIND_H_
* \return A schedule for the given ops.
*/
inline Schedule schedule_dense(const Target& target, const Array<Tensor>& outs) {
- if (target->id->name == "cuda" && target->GetLibs().count("cublas")) {
+ if (target->kind->name == "cuda" && target->GetLibs().count("cublas")) {
return topi::generic::schedule_extern(target, outs);
}
if (out_stage->op.as<ComputeOpNode>()->axis.size() > 0) {
all_reduce = false;
num_thread = 32;
- if (target->id->name == "opencl") {
+ if (target->kind->name == "opencl") {
// Without this, CL_INVALID_WORK_GROUP_SIZE occurs with python tests.
// Don't know why.
num_thread = 16;
* \return A schedule for the given ops.
*/
inline Schedule schedule_dense(const Target& target, const Array<Tensor>& outs) {
- if (target->id->name == "rocm" && target->GetLibs().count("rocblas")) {
+ if (target->kind->name == "rocm" && target->GetLibs().count("rocblas")) {
return topi::generic::schedule_extern(target, outs);
}
continue
if workload_key and inp.task.workload_key != workload_key:
continue
- if target and inp.task.target.id.name != target.id.name:
+ if target and inp.task.target.kind.name != target.kind.name:
continue
costs = [v.value for v in res.costs]
device = tgt.attrs.get("device", "")
if device != "":
possible_names.append(_alias(device))
- possible_names.append(tgt.id.name)
+ possible_names.append(tgt.kind.name)
all_packages = list(PACKAGE_VERSION.keys())
for name in possible_names:
"""
# config setup
pass_ctx = PassContext.current()
- instrument_bound_checkers = bool(pass_ctx.config.get("tir.instrument_bound_checkers", False))
- disable_vectorize = bool(pass_ctx.config.get("tir.disable_vectorize", False))
+ instrument_bound_checkers = bool(pass_ctx.config.get(
+ "tir.instrument_bound_checkers", False))
+ disable_vectorize = bool(pass_ctx.config.get(
+ "tir.disable_vectorize", False))
add_lower_pass = pass_ctx.config.get("tir.add_lower_pass", [])
lower_phase0 = [x[1] for x in add_lower_pass if x[0] == 0]
"""
target = _target.create(target)
target_host = _target.create(target_host)
- device_type = ndarray.context(target.id.name, 0).device_type
+ device_type = ndarray.context(target.kind.name, 0).device_type
mod_mixed = input_mod
- mod_mixed = tvm.tir.transform.Apply(lambda f: f.with_attr("target", target))(mod_mixed)
+ mod_mixed = tvm.tir.transform.Apply(
+ lambda f: f.with_attr("target", target))(mod_mixed)
opt_mixed = [tvm.tir.transform.VerifyMemory()]
if len(mod_mixed.functions) == 1:
- opt_mixed += [tvm.tir.transform.Apply(lambda f: f.with_attr("tir.is_entry_func", True))]
+ opt_mixed += [tvm.tir.transform.Apply(
+ lambda f: f.with_attr("tir.is_entry_func", True))]
if PassContext.current().config.get("tir.detect_global_barrier", False):
opt_mixed += [tvm.tir.transform.ThreadSync("global")]
tvm.tir.transform.SplitHostDevice()]
mod_mixed = tvm.transform.Sequential(opt_mixed)(mod_mixed)
-
# device optimizations
opt_device = tvm.transform.Sequential(
[tvm.tir.transform.Filter(
"Specified target %s, but cannot find device code, did you do "
"bind?" % target)
- rt_mod_dev = codegen.build_module(mod_dev, target) if len(mod_dev.functions) != 0 else None
+ rt_mod_dev = codegen.build_module(mod_dev, target) if len(
+ mod_dev.functions) != 0 else None
return mod_host, rt_mod_dev
elif isinstance(inputs, tvm.IRModule):
input_mod = inputs
elif not isinstance(inputs, (dict, container.Map)):
- raise ValueError("inputs must be Schedule, IRModule or dict of target to IRModule")
+ raise ValueError(
+ "inputs must be Schedule, IRModule or dict of target to IRModule")
if not isinstance(inputs, (dict, container.Map)):
target = _target.Target.current() if target is None else target
if not target_host:
for tar, _ in target_input_mod.items():
tar = _target.create(tar)
- device_type = ndarray.context(tar.id.name, 0).device_type
+ device_type = ndarray.context(tar.kind.name, 0).device_type
if device_type == ndarray.cpu(0).device_type:
target_host = tar
break
wrap_compute_softmax(topi.nn.softmax),
wrap_topi_schedule(topi.cuda.schedule_softmax),
name="softmax.cuda")
- if target.id.name == "cuda" and "cudnn" in target.libs:
+ if target.kind.name == "cuda" and "cudnn" in target.libs:
strategy.add_implementation(
wrap_compute_softmax(topi.cuda.softmax_cudnn),
wrap_topi_schedule(topi.cuda.schedule_softmax_cudnn),
dilation_h, dilation_w,
pre_flag=False)
if judge_winograd_shape:
- if target.id.name == "cuda" and \
+ if target.kind.name == "cuda" and \
nvcc.have_tensorcore(tvm.gpu(0).compute_version) and \
judge_winograd_tensorcore:
strategy.add_implementation(
topi.cuda.schedule_conv2d_nhwc_winograd_direct),
name="conv2d_nhwc_winograd_direct.cuda",
plevel=5)
- if target.id.name == "cuda":
+ if target.kind.name == "cuda":
if nvcc.have_tensorcore(tvm.gpu(0).compute_version):
if (N % 16 == 0 and CI % 16 == 0 and CO % 16 == 0) or \
(N % 8 == 0 and CI % 16 == 0 and CO % 32 == 0) or \
else:
raise RuntimeError("Unsupported conv2d layout {} for CUDA".format(layout))
# add cudnn implementation
- if target.id.name == "cuda" and "cudnn" in target.libs:
+ if target.kind.name == "cuda" and "cudnn" in target.libs:
if layout in ["NCHW", "NHWC"] and padding[0] == padding[2] and \
padding[1] == padding[3]:
strategy.add_implementation(
else: # group_conv2d
# add cudnn implementation, if any
cudnn_impl = False
- if target.id.name == "cuda" and "cudnn" in target.libs:
+ if target.kind.name == "cuda" and "cudnn" in target.libs:
if layout in ["NCHW", "NHWC"] and padding[0] == padding[2] and \
padding[1] == padding[3]:
strategy.add_implementation(
padding, stride_h, stride_w,
dilation_h, dilation_w,
pre_flag=True)
- if target.id.name == "cuda" and \
+ if target.kind.name == "cuda" and \
nvcc.have_tensorcore(tvm.gpu(0).compute_version) and \
judge_winograd_tensorcore:
strategy.add_implementation(
plevel=10)
N, _, _, _, _ = get_const_tuple(data.shape)
_, _, _, CI, CO = get_const_tuple(kernel.shape)
- if target.id.name == "cuda":
+ if target.kind.name == "cuda":
if nvcc.have_tensorcore(tvm.gpu(0).compute_version):
if (N % 16 == 0 and CI % 16 == 0 and CO % 16 == 0) or \
(N % 8 == 0 and CI % 16 == 0 and CO % 32 == 0) or \
name="conv3d_ndhwc_tensorcore.cuda",
plevel=20)
- if target.id.name == "cuda" and "cudnn" in target.libs:
+ if target.kind.name == "cuda" and "cudnn" in target.libs:
strategy.add_implementation(wrap_compute_conv3d(topi.cuda.conv3d_cudnn, True),
wrap_topi_schedule(topi.cuda.schedule_conv3d_cudnn),
name="conv3d_cudnn.cuda",
wrap_topi_schedule(topi.cuda.schedule_dense_large_batch),
name="dense_large_batch.cuda",
plevel=5)
- if target.id.name == "cuda":
+ if target.kind.name == "cuda":
if nvcc.have_tensorcore(tvm.gpu(0).compute_version):
if(i % 16 == 0 and b % 16 == 0 and o % 16 == 0) \
or (i % 16 == 0 and b % 8 == 0 and o % 32 == 0) \
wrap_topi_schedule(topi.cuda.schedule_dense_tensorcore),
name="dense_tensorcore.cuda",
plevel=20)
- if target.id.name == "cuda" and "cublas" in target.libs:
+ if target.kind.name == "cuda" and "cublas" in target.libs:
strategy.add_implementation(
wrap_compute_dense(topi.cuda.dense_cublas),
wrap_topi_schedule(topi.cuda.schedule_dense_cublas),
wrap_topi_schedule(topi.cuda.schedule_batch_matmul),
name="batch_matmul.cuda",
plevel=10)
- if target.id.name == "cuda" and "cublas" in target.libs:
+ if target.kind.name == "cuda" and "cublas" in target.libs:
strategy.add_implementation(
wrap_compute_batch_matmul(topi.cuda.batch_matmul_cublas),
wrap_topi_schedule(topi.generic.schedule_extern),
wrap_compute_dense(topi.rocm.dense),
wrap_topi_schedule(topi.rocm.schedule_dense),
name="dense.rocm")
- if target.id.name == "rocm" and "rocblas" in target.libs:
+ if target.kind.name == "rocm" and "rocblas" in target.libs:
assert out_type.dtype == inputs[0].dtype, "Mixed precision not supported."
strategy.add_implementation(
wrap_compute_dense(topi.rocm.dense_rocblas),
if tvm.target.Target.current():
target = tvm.target.Target.current()
- ctx = tvm.context(target.id.name)
+ ctx = tvm.context(target.kind.name)
else:
target = 'llvm'
ctx = tvm.context(target)
# under the License.
"""Target description and codgen module.
-TVM's target string is in fomat ``<target_id> [-option=value]...``.
+TVM's target string is in fomat ``<target_kind> [-option=value]...``.
Note
----
@tvm._ffi.register_object
-class TargetId(Object):
- """Id of a compilation target
+class TargetKind(Object):
+ """Kind of a compilation target
"""
- :py:func:`tvm.target.mali` create Mali target
- :py:func:`tvm.target.intel_graphics` create Intel Graphics target
"""
+
def __enter__(self):
_ffi_api.EnterTargetScope(self)
return self
options : str or list of str
Additional options
"""
- opts = ["-device=intel_graphics", "-model=%s" % model, "-thread_warp_size=16"]
+ opts = ["-device=intel_graphics", "-model=%s" %
+ model, "-thread_warp_size=16"]
opts = _merge_opts(opts, options)
return _ffi_api.TargetCreate("opencl", *opts)
i = sim_args.index('hvx_length') + len('hvx_length') + 1
sim_hvx = sim_args[i:i+3]
if sim_hvx != str(codegen_hvx):
- print('WARNING: sim hvx {} and codegen hvx {} mismatch!' \
+ print('WARNING: sim hvx {} and codegen hvx {} mismatch!'
.format(sim_hvx, codegen_hvx))
elif codegen_hvx != 0:
# If --hvx_length was not given, add it if HVX is enabled
# Parse options into correct order
cpu_attr = {x: str(m.groupdict()[x] or '') for x in m.groupdict()}
sim_args = cpu_attr['base_version'] + \
- cpu_attr['sub_version'] + \
- cpu_attr['l2_size'] + \
- cpu_attr['rev'] + ' ' + \
- cpu_attr['pre'] + cpu_attr['post']
+ cpu_attr['sub_version'] + \
+ cpu_attr['l2_size'] + \
+ cpu_attr['rev'] + ' ' + \
+ cpu_attr['pre'] + cpu_attr['post']
return sim_cpu + ' ' + validate_hvx_length(hvx, sim_args)
cfg.define_split("tile_k", k, num_outputs=2)
cfg.define_knob("auto_unroll_max_step", [8, 16, 32, 64])
target = tvm.target.Target.current()
- if target.id.name in ['nvptx', 'rocm']:
+ if target.kind.name in ['nvptx', 'rocm']:
# llvm-based backends cannot do non-explicit unrolling
cfg.define_knob("unroll_explicit", [1])
else:
cfg.define_knob("auto_unroll_max_step", [64, 512, 1500])
target = tvm.target.Target.current()
- if target.id.name in ['nvptx', 'rocm']:
+ if target.kind.name in ['nvptx', 'rocm']:
cfg.define_knob("unroll_explicit", [1])
else:
cfg.define_knob("unroll_explicit", [0, 1])
cfg.define_knob("auto_unroll_max_step", [64, 512, 1500])
target = tvm.target.Target.current()
- if target.id.name in ['nvptx', 'rocm']:
+ if target.kind.name in ['nvptx', 'rocm']:
cfg.define_knob("unroll_explicit", [1])
else:
cfg.define_knob("unroll_explicit", [0, 1])
cfg.define_knob("auto_unroll_max_step", [64, 512, 1500])
target = tvm.target.Target.current()
- if target.id.name in ['nvptx', 'rocm']:
+ if target.kind.name in ['nvptx', 'rocm']:
cfg.define_knob("unroll_explicit", [1])
else:
cfg.define_knob("unroll_explicit", [0, 1])
cfg.define_knob("auto_unroll_max_step", [0, 512, 1500])
target = tvm.target.Target.current()
- if target.id.name in ['nvptx', 'rocm']:
+ if target.kind.name in ['nvptx', 'rocm']:
cfg.define_knob("unroll_explicit", [1])
else:
cfg.define_knob("unroll_explicit", [0, 1])
# fallback support
if cfg.is_fallback:
ref_log = autotvm.tophub.load_reference_log(
- target.id.name, target.model, 'conv2d_nchw.cuda')
+ target.kind.name, target.model, 'conv2d_nchw.cuda')
cfg.fallback_with_reference_log(ref_log)
##### space definition end #####
target = tvm.target.Target.current()
if cfg.is_fallback:
ref_log = autotvm.tophub.load_reference_log(
- target.id.name, target.model, 'conv2d_nhwc.cuda')
+ target.kind.name, target.model, 'conv2d_nhwc.cuda')
cfg.fallback_with_reference_log(ref_log)
tile_n = cfg["tile_n"].val
target = tvm.target.Target.current()
if cfg.is_fallback:
ref_log = autotvm.tophub.load_reference_log(
- target.id.name, target.model, 'conv2d_nhwc_tensorcore.cuda')
+ target.kind.name, target.model, 'conv2d_nhwc_tensorcore.cuda')
cfg.fallback_with_reference_log(ref_log)
block_row_warps = cfg["block_row_warps"].val
cfg.define_knob("auto_unroll_max_step", [64, 512, 1500])
target = tvm.target.Target.current()
- if target.id.name in ['nvptx', 'rocm']:
+ if target.kind.name in ['nvptx', 'rocm']:
cfg.define_knob("unroll_explicit", [1])
else:
cfg.define_knob("unroll_explicit", [0, 1])
cfg.define_split("tile_rc", rc, num_outputs=2)
cfg.define_knob("auto_unroll_max_step", [0, 128, 1500])
target = tvm.target.Target.current()
- if target.id.name in ['nvptx', 'rocm']:
+ if target.kind.name in ['nvptx', 'rocm']:
cfg.define_knob("unroll_explicit", [1])
else:
cfg.define_knob("unroll_explicit", [0, 1])
cfg.define_knob("auto_unroll_max_step", [0, 512, 1500])
target = tvm.target.Target.current()
- if target.id.name in ['nvptx', 'rocm']:
+ if target.kind.name in ['nvptx', 'rocm']:
cfg.define_knob("unroll_explicit", [1])
else:
cfg.define_knob("unroll_explicit", [0, 1])
# fallback support
if cfg.is_fallback:
ref_log = autotvm.tophub.load_reference_log(
- target.id.name, target.model, workload_name)
+ target.kind.name, target.model, workload_name)
cfg.fallback_with_reference_log(ref_log)
##### space definition end #####
target = tvm.target.Target.current()
if cfg.is_fallback:
ref_log = autotvm.tophub.load_reference_log(
- target.id.name, target.model, 'conv3d_ndhwc_tensorcore.cuda')
+ target.kind.name, target.model, 'conv3d_ndhwc_tensorcore.cuda')
cfg.fallback_with_reference_log(ref_log)
block_row_warps = cfg["block_row_warps"].val
cfg.define_split("tile_rc", rc, num_outputs=2)
cfg.define_knob("auto_unroll_max_step", [0, 128, 1500])
target = tvm.target.Target.current()
- if target.id.name in ['nvptx', 'rocm']:
+ if target.kind.name in ['nvptx', 'rocm']:
cfg.define_knob("unroll_explicit", [1])
else:
cfg.define_knob("unroll_explicit", [0, 1])
cfg.define_split("tile_rz", rz, num_outputs=2)
cfg.define_knob("auto_unroll_max_step", [0, 128, 1500])
target = tvm.target.Target.current()
- if target.id.name in ['nvptx', 'rocm']:
+ if target.kind.name in ['nvptx', 'rocm']:
cfg.define_knob("unroll_explicit", [1])
else:
cfg.define_knob("unroll_explicit", [0, 1])
cfg.define_knob("auto_unroll_max_step", [0, 512, 1500])
target = tvm.target.Target.current()
- if target.id.name in ['nvptx', 'rocm']:
+ if target.kind.name in ['nvptx', 'rocm']:
cfg.define_knob("unroll_explicit", [1])
else:
cfg.define_knob("unroll_explicit", [0, 1])
s[load].bind(tx, te.thread_axis("threadIdx.x"))
# unroll
- s[output].pragma(kernel_scope, 'auto_unroll_max_step', cfg['auto_unroll_max_step'].val)
- s[output].pragma(kernel_scope, 'unroll_explicit', cfg['unroll_explicit'].val)
+ s[output].pragma(kernel_scope, 'auto_unroll_max_step',
+ cfg['auto_unroll_max_step'].val)
+ s[output].pragma(kernel_scope, 'unroll_explicit',
+ cfg['unroll_explicit'].val)
@autotvm.register_topi_schedule("correlation_nchw.cuda")
cfg.define_knob("auto_unroll_max_step", [0, 512, 1500])
target = tvm.target.Target.current()
- if target.id.name in ['nvptx', 'rocm']:
+ if target.kind.name in ['nvptx', 'rocm']:
cfg.define_knob("unroll_explicit", [1])
else:
cfg.define_knob("unroll_explicit", [0, 1])
target = tvm.target.Target.current()
if cfg.is_fallback:
ref_log = autotvm.tophub.load_reference_log(
- target.id.name, target.model, 'dense_tensorcore.cuda')
+ target.kind.name, target.model, 'dense_tensorcore.cuda')
cfg.fallback_with_reference_log(ref_log)
# Deal with op fusion, such as bias and relu
cfg.define_knob("auto_unroll_max_step", [0, 256, 1500])
target = tvm.target.Target.current()
- if target.id.name in ['nvptx', 'rocm']:
+ if target.kind.name in ['nvptx', 'rocm']:
cfg.define_knob("unroll_explicit", [1])
else:
cfg.define_knob("unroll_explicit", [0, 1])
# fallback support
if cfg.is_fallback:
ref_log = autotvm.tophub.load_reference_log(
- target.id.name, target.model, 'depthwise_conv2d_nchw.cuda')
+ target.kind.name, target.model, 'depthwise_conv2d_nchw.cuda')
cfg.fallback_with_reference_log(ref_log)
# TODO(lmzheng): A bug here, set unroll_explicit to False as workaround
cfg['unroll_explicit'].val = 0
# num_thread here could be 728, it is larger than cuda.max_num_threads
num_thread = tvm.arith.Analyzer().simplify(temp.shape[3]).value
target = tvm.target.Target.current()
- if target and (target.id.name not in ["cuda", "nvptx"]):
+ if target and (target.kind.name not in ["cuda", "nvptx"]):
num_thread = target.max_num_threads
xoc, xic = s[Output].split(c, factor=num_thread)
s[Output].reorder(xoc, b, h, w, xic)
cfg.define_knob("auto_unroll_max_step", [0, 512, 1500])
target = tvm.target.Target.current()
- if target.id.name in ['nvptx', 'rocm']:
+ if target.kind.name in ['nvptx', 'rocm']:
cfg.define_knob("unroll_explicit", [1])
else:
cfg.define_knob("unroll_explicit", [0, 1])
all_reduce = False
num_thread = 32
target = tvm.target.Target.current()
- if target and target.id.name == "opencl":
+ if target and target.kind.name == "opencl":
# without it, CL_INVALID_WORK_GROUP_SIZE occurred when running test_topi_reduce.py
# don't know why
num_thread = 16
#
# TODO(tvm-team) Fix nvptx codegen or deprecate nvptx backend.
def sched_warp_softmax():
- if tgt.id.name == "nvptx" or tgt.id.name == "rocm":
+ if tgt.kind.name == "nvptx" or tgt.kind.name == "rocm":
return softmax.dtype == "float32" or softmax.dtype == "int32"
- if tgt.id.name != "cuda":
+ if tgt.kind.name != "cuda":
# this is used as the gpu schedule for other arches which may not have warp reductions
return False
return True
The computation schedule for reorg.
"""
target = tvm.target.Target.current(allow_none=False)
- cpp_target = cpp.TEST_create_target(target.id.name)
+ cpp_target = cpp.TEST_create_target(target.kind.name)
return cpp.cuda.schedule_injective(cpp_target, outs)
def schedule_nms(outs):
"""Default schedule for llvm."""
target = tvm.target.Target.current(allow_none=False)
outs = [outs] if isinstance(outs, te.tensor.Tensor) else outs
- if target.id.name not in ("llvm", "c"):
+ if target.kind.name not in ("llvm", "c"):
raise RuntimeError("schedule not registered for '%s'" % target)
s = te.create_schedule([x.op for x in outs])
if auto_inline:
The computation schedule for the op.
"""
target = tvm.target.Target.current(allow_none=False)
- if target.id.name != "llvm":
+ if target.kind.name != "llvm":
raise RuntimeError("schedule_injective not registered for '%s'" % target)
outs = [outs] if isinstance(outs, te.tensor.Tensor) else outs
x = outs[0]
The computation schedule for the op.
"""
target = tvm.target.Target.current(allow_none=False)
- cpp_target = cpp.TEST_create_target(target.id.name)
+ cpp_target = cpp.TEST_create_target(target.kind.name)
return cpp.generic.default_schedule(cpp_target, outs, False)
def schedule_get_valid_counts(outs):
cfg.define_knob("auto_unroll_max_step", [0, 256, 1500])
target = tvm.target.Target.current()
- if target.id.name in ['nvptx', 'rocm']:
+ if target.kind.name in ['nvptx', 'rocm']:
cfg.define_knob("unroll_explicit", [1])
else:
cfg.define_knob("unroll_explicit", [0, 1])
# fallback support
if cfg.is_fallback:
ref_log = autotvm.tophub.load_reference_log(
- target.id.name, target.model, 'depthwise_conv2d_nchw.intel_graphics')
+ target.kind.name, target.model, 'depthwise_conv2d_nchw.intel_graphics')
cfg.fallback_with_reference_log(ref_log)
cfg['unroll_explicit'].val = 0
##### space definition end #####
# num_thread here could be 728, it is larger than cuda.max_num_threads
num_thread = tvm.arith.Analyzer().simplify(temp.shape[3]).value
target = tvm.target.Target.current()
- if target and (target.id.name not in ["cuda", "nvptx"]):
+ if target and (target.kind.name not in ["cuda", "nvptx"]):
num_thread = target.max_num_threads
xoc, xic = s[Output].split(c, factor=num_thread)
s[Output].reorder(xoc, b, h, w, xic)
HardwareParams HardwareParamsNode::GetDefaultHardwareParams(const Target& target,
const Target& target_host) {
- if (target->id->name == "llvm") {
+ if (target->kind->name == "llvm") {
return HardwareParams(tvm::runtime::threading::MaxConcurrency(), 64, 64);
} else {
LOG(FATAL) << "No default hardware parameters for target: " << target;
/*! \return The default host target for a given device target */
Target DefaultTargetHost(Target target) {
- if (target.defined() && target->id->device_type == kDLCPU) {
+ if (target.defined() && target->kind->device_type == kDLCPU) {
return target;
} else {
if (LLVMEnabled()) {
<< " but cannot find device code. Did you forget to bind?";
}
- if (target->id->device_type == kDLCPU && target_host == target) {
+ if (target->kind->device_type == kDLCPU && target_host == target) {
CHECK(mdevice->functions.empty()) << "No device code should be generated when target "
<< "and host_target are both llvm target."
<< "\n";
Target target_host_val = target_host;
if (!target_host.defined()) {
for (const auto& it : inputs) {
- if (it.first->id->device_type == kDLCPU || it.first->id->device_type == kDLMicroDev) {
+ if (it.first->kind->device_type == kDLCPU || it.first->kind->device_type == kDLMicroDev) {
target_host_val = it.first;
break;
}
if (!target_host.defined())
target_host = (pf != nullptr) ? target::llvm() : target::stackvm();
- if (target_host.defined() && target_host->id->name == "llvm") {
+ if (target_host.defined() && target_host->kind->name == "llvm") {
// If we can decide the target is LLVM, we then create an empty LLVM module.
ret_.mod = (*pf)(target_host->str(), "empty_module");
} else {
Target target_host = target_host_;
if (!target_host_.defined()) {
for (const auto& it : targets_) {
- if (it.second->id->device_type == kDLCPU) {
+ if (it.second->kind->device_type == kDLCPU) {
target_host = it.second;
break;
}
mod = tir::transform::SkipAssert()(mod);
}
std::string build_f_name;
- if (target->id->name == "micro_dev") {
+ if (target->kind->name == "micro_dev") {
build_f_name = "target.build.c";
} else {
- build_f_name = "target.build." + target->id->name;
+ build_f_name = "target.build." + target->kind->name;
}
// the build function.
const PackedFunc* bf = runtime::Registry::Get(build_f_name);
#include <tvm/node/repr_printer.h>
#include <tvm/runtime/registry.h>
#include <tvm/target/target.h>
-#include <tvm/target/target_id.h>
+#include <tvm/target/target_kind.h>
#include <tvm/tir/expr.h>
#include <algorithm>
using runtime::TVMRetValue;
Target Target::CreateTarget(const std::string& name, const std::vector<std::string>& options) {
- TargetId id = TargetId::Get(name);
+ TargetKind kind = TargetKind::Get(name);
ObjectPtr<TargetNode> target = make_object<TargetNode>();
- target->id = id;
+ target->kind = kind;
// tag is always empty
target->tag = "";
// parse attrs
- target->attrs = id->ParseAttrsFromRaw(options);
+ target->attrs = kind->ParseAttrsFromRaw(options);
String device_name = target->GetAttr<String>("device", "").value();
// set up keys
{
keys.push_back(device_name);
}
// add default keys
- for (const auto& key : target->id->default_keys) {
+ for (const auto& key : target->kind->default_keys) {
keys.push_back(key);
}
// de-duplicate keys
const std::string& TargetNode::str() const {
if (str_repr_.empty()) {
std::ostringstream os;
- os << id->name;
+ os << kind->name;
if (!this->keys.empty()) {
os << " -keys=";
bool is_first = true;
os << s;
}
}
- if (Optional<String> attrs_str = id->StringifyAttrsToRaw(attrs)) {
+ if (Optional<String> attrs_str = kind->StringifyAttrsToRaw(attrs)) {
os << ' ' << attrs_str.value();
}
str_repr_ = os.str();
*/
/*!
- * \file src/target/target_id.cc
- * \brief Target id registry
+ * \file src/target/target_kind.cc
+ * \brief Target kind registry
*/
-#include <tvm/target/target_id.h>
+#include <tvm/target/target_kind.h>
#include <algorithm>
namespace tvm {
-TVM_REGISTER_NODE_TYPE(TargetIdNode);
+TVM_REGISTER_NODE_TYPE(TargetKindNode);
TVM_STATIC_IR_FUNCTOR(ReprPrinter, vtable)
- .set_dispatch<TargetIdNode>([](const ObjectRef& node, ReprPrinter* p) {
- auto* op = static_cast<const TargetIdNode*>(node.get());
+ .set_dispatch<TargetKindNode>([](const ObjectRef& node, ReprPrinter* p) {
+ auto* op = static_cast<const TargetKindNode*>(node.get());
p->stream << op->name;
});
-using TargetIdRegistry = AttrRegistry<TargetIdRegEntry, TargetId>;
+using TargetKindRegistry = AttrRegistry<TargetKindRegEntry, TargetKind>;
-TargetIdRegEntry& TargetIdRegEntry::RegisterOrGet(const String& target_id_name) {
- return TargetIdRegistry::Global()->RegisterOrGet(target_id_name);
+TargetKindRegEntry& TargetKindRegEntry::RegisterOrGet(const String& target_kind_name) {
+ return TargetKindRegistry::Global()->RegisterOrGet(target_kind_name);
}
-void TargetIdRegEntry::UpdateAttr(const String& key, TVMRetValue value, int plevel) {
- TargetIdRegistry::Global()->UpdateAttr(key, id_, value, plevel);
+void TargetKindRegEntry::UpdateAttr(const String& key, TVMRetValue value, int plevel) {
+ TargetKindRegistry::Global()->UpdateAttr(key, kind_, value, plevel);
}
-const AttrRegistryMapContainerMap<TargetId>& TargetId::GetAttrMapContainer(
+const AttrRegistryMapContainerMap<TargetKind>& TargetKind::GetAttrMapContainer(
const String& attr_name) {
- return TargetIdRegistry::Global()->GetAttrMap(attr_name);
+ return TargetKindRegistry::Global()->GetAttrMap(attr_name);
}
-const TargetId& TargetId::Get(const String& target_id_name) {
- const TargetIdRegEntry* reg = TargetIdRegistry::Global()->Get(target_id_name);
- CHECK(reg != nullptr) << "ValueError: TargetId \"" << target_id_name << "\" is not registered";
- return reg->id_;
+const TargetKind& TargetKind::Get(const String& target_kind_name) {
+ const TargetKindRegEntry* reg = TargetKindRegistry::Global()->Get(target_kind_name);
+ CHECK(reg != nullptr) << "ValueError: TargetKind \"" << target_kind_name
+ << "\" is not registered";
+ return reg->kind_;
}
-void TargetIdNode::VerifyTypeInfo(const ObjectRef& obj,
- const TargetIdNode::ValueTypeInfo& info) const {
+void TargetKindNode::VerifyTypeInfo(const ObjectRef& obj,
+ const TargetKindNode::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 "
}
}
-void TargetIdNode::ValidateSchema(const Map<String, ObjectRef>& config) const {
- const String kTargetId = "id";
+void TargetKindNode::ValidateSchema(const Map<String, ObjectRef>& config) const {
+ const String kTargetKind = "kind";
for (const auto& kv : config) {
const String& name = kv.first;
const ObjectRef& obj = kv.second;
- if (name == kTargetId) {
+ if (name == kTargetKind) {
CHECK(obj->IsInstance<StringObj>())
- << "AttributeError: \"id\" is not a string, but its type is \"" << obj->GetTypeKey()
+ << "AttributeError: \"kind\" is not a string, but its type is \"" << obj->GetTypeKey()
<< "\"";
CHECK(Downcast<String>(obj) == this->name)
- << "AttributeError: \"id\" = \"" << obj << "\" is inconsistent with TargetId \""
+ << "AttributeError: \"kind\" = \"" << obj << "\" is inconsistent with TargetKind \""
<< this->name << "\"";
continue;
}
try {
VerifyTypeInfo(obj, info);
} catch (const tvm::Error& e) {
- LOG(FATAL) << "AttributeError: Schema validation failed for TargetId \"" << this->name
+ LOG(FATAL) << "AttributeError: Schema validation failed for TargetKind \"" << this->name
<< "\", details:\n"
<< e.what() << "\n"
<< "The config is:\n"
}
}
-inline String GetId(const Map<String, ObjectRef>& target, const char* name) {
- const String kTargetId = "id";
- CHECK(target.count(kTargetId)) << "AttributeError: \"id\" does not exist in \"" << name << "\"\n"
- << name << " = " << target;
- const ObjectRef& obj = target[kTargetId];
- CHECK(obj->IsInstance<StringObj>()) << "AttributeError: \"id\" is not a string in \"" << name
+inline String GetKind(const Map<String, ObjectRef>& target, const char* name) {
+ const String kTargetKind = "kind";
+ CHECK(target.count(kTargetKind))
+ << "AttributeError: \"kind\" does not exist in \"" << name << "\"\n"
+ << name << " = " << target;
+ const ObjectRef& obj = target[kTargetKind];
+ CHECK(obj->IsInstance<StringObj>()) << "AttributeError: \"kind\" is not a string in \"" << name
<< "\", but its type is \"" << obj->GetTypeKey() << "\"\n"
<< name << " = \"" << target << '"';
return Downcast<String>(obj);
const String kTargetHost = "target_host";
Map<String, ObjectRef> target = config;
Map<String, ObjectRef> target_host;
- String target_id = GetId(target, "target");
- String target_host_id;
+ String target_kind = GetKind(target, "target");
+ String target_host_kind;
if (config.count(kTargetHost)) {
target.erase(kTargetHost);
target_host = Downcast<Map<String, ObjectRef>>(config[kTargetHost]);
- target_host_id = GetId(target_host, "target_host");
+ target_host_kind = GetKind(target_host, "target_host");
}
- TargetId::Get(target_id)->ValidateSchema(target);
+ TargetKind::Get(target_kind)->ValidateSchema(target);
if (!target_host.empty()) {
- TargetId::Get(target_host_id)->ValidateSchema(target_host);
+ TargetKind::Get(target_host_kind)->ValidateSchema(target_host);
}
} catch (const tvm::Error& e) {
LOG(FATAL) << "AttributeError: schedule validation fails:\n"
return String(os.str());
}
-Map<String, ObjectRef> TargetIdNode::ParseAttrsFromRaw(
+Map<String, ObjectRef> TargetKindNode::ParseAttrsFromRaw(
const std::vector<std::string>& options) const {
std::unordered_map<String, ObjectRef> attrs;
for (size_t iter = 0, end = options.size(); iter < end;) {
return attrs;
}
-Optional<String> TargetIdNode::StringifyAttrsToRaw(const Map<String, ObjectRef>& attrs) const {
+Optional<String> TargetKindNode::StringifyAttrsToRaw(const Map<String, ObjectRef>& attrs) const {
std::ostringstream os;
std::vector<String> keys;
for (const auto& kv : attrs) {
// TODO(@junrushao1994): remove some redundant attributes
-TVM_REGISTER_TARGET_ID("llvm")
+TVM_REGISTER_TARGET_KIND("llvm")
.add_attr_option<Array<String>>("keys")
.add_attr_option<Array<String>>("libs")
.add_attr_option<String>("device")
.set_default_keys({"cpu"})
.set_device_type(kDLCPU);
-TVM_REGISTER_TARGET_ID("c")
+TVM_REGISTER_TARGET_KIND("c")
.add_attr_option<Array<String>>("keys")
.add_attr_option<Array<String>>("libs")
.add_attr_option<String>("device")
.set_default_keys({"cpu"})
.set_device_type(kDLCPU);
-TVM_REGISTER_TARGET_ID("micro_dev")
+TVM_REGISTER_TARGET_KIND("micro_dev")
.add_attr_option<Array<String>>("keys")
.add_attr_option<Array<String>>("libs")
.add_attr_option<String>("device")
.set_default_keys({"micro_dev"})
.set_device_type(kDLMicroDev);
-TVM_REGISTER_TARGET_ID("cuda")
+TVM_REGISTER_TARGET_KIND("cuda")
.add_attr_option<Array<String>>("keys")
.add_attr_option<Array<String>>("libs")
.add_attr_option<String>("device")
.set_default_keys({"cuda", "gpu"})
.set_device_type(kDLGPU);
-TVM_REGISTER_TARGET_ID("nvptx")
+TVM_REGISTER_TARGET_KIND("nvptx")
.add_attr_option<Array<String>>("keys")
.add_attr_option<Array<String>>("libs")
.add_attr_option<String>("device")
.set_default_keys({"cuda", "gpu"})
.set_device_type(kDLGPU);
-TVM_REGISTER_TARGET_ID("rocm")
+TVM_REGISTER_TARGET_KIND("rocm")
.add_attr_option<Array<String>>("keys")
.add_attr_option<Array<String>>("libs")
.add_attr_option<String>("device")
.set_default_keys({"rocm", "gpu"})
.set_device_type(kDLROCM);
-TVM_REGISTER_TARGET_ID("opencl")
+TVM_REGISTER_TARGET_KIND("opencl")
.add_attr_option<Array<String>>("keys")
.add_attr_option<Array<String>>("libs")
.add_attr_option<String>("device")
.set_default_keys({"opencl", "gpu"})
.set_device_type(kDLOpenCL);
-TVM_REGISTER_TARGET_ID("metal")
+TVM_REGISTER_TARGET_KIND("metal")
.add_attr_option<Array<String>>("keys")
.add_attr_option<Array<String>>("libs")
.add_attr_option<String>("device")
.set_default_keys({"metal", "gpu"})
.set_device_type(kDLMetal);
-TVM_REGISTER_TARGET_ID("vulkan")
+TVM_REGISTER_TARGET_KIND("vulkan")
.add_attr_option<Array<String>>("keys")
.add_attr_option<Array<String>>("libs")
.add_attr_option<String>("device")
.set_default_keys({"vulkan", "gpu"})
.set_device_type(kDLVulkan);
-TVM_REGISTER_TARGET_ID("webgpu")
+TVM_REGISTER_TARGET_KIND("webgpu")
.add_attr_option<Array<String>>("keys")
.add_attr_option<Array<String>>("libs")
.add_attr_option<String>("device")
.set_default_keys({"webgpu", "gpu"})
.set_device_type(kDLWebGPU);
-TVM_REGISTER_TARGET_ID("sdaccel")
+TVM_REGISTER_TARGET_KIND("sdaccel")
.add_attr_option<Array<String>>("keys")
.add_attr_option<Array<String>>("libs")
.add_attr_option<String>("device")
.set_default_keys({"sdaccel", "hls"})
.set_device_type(kDLOpenCL);
-TVM_REGISTER_TARGET_ID("aocl")
+TVM_REGISTER_TARGET_KIND("aocl")
.add_attr_option<Array<String>>("keys")
.add_attr_option<Array<String>>("libs")
.add_attr_option<String>("device")
.set_default_keys({"aocl", "hls"})
.set_device_type(kDLAOCL);
-TVM_REGISTER_TARGET_ID("aocl_sw_emu")
+TVM_REGISTER_TARGET_KIND("aocl_sw_emu")
.add_attr_option<Array<String>>("keys")
.add_attr_option<Array<String>>("libs")
.add_attr_option<String>("device")
.set_default_keys({"aocl", "hls"})
.set_device_type(kDLAOCL);
-TVM_REGISTER_TARGET_ID("hexagon")
+TVM_REGISTER_TARGET_KIND("hexagon")
.add_attr_option<Array<String>>("keys")
.add_attr_option<Array<String>>("libs")
.add_attr_option<String>("device")
.set_default_keys({"hexagon"})
.set_device_type(kDLHexagon);
-TVM_REGISTER_TARGET_ID("stackvm")
+TVM_REGISTER_TARGET_KIND("stackvm")
.add_attr_option<Array<String>>("keys")
.add_attr_option<Array<String>>("libs")
.add_attr_option<String>("device")
.add_attr_option<Bool>("system-lib")
.set_device_type(kDLCPU);
-TVM_REGISTER_TARGET_ID("ext_dev")
+TVM_REGISTER_TARGET_KIND("ext_dev")
.add_attr_option<Array<String>>("keys")
.add_attr_option<Array<String>>("libs")
.add_attr_option<String>("device")
.add_attr_option<Bool>("system-lib")
.set_device_type(kDLExtDev);
-TVM_REGISTER_TARGET_ID("hybrid")
+TVM_REGISTER_TARGET_KIND("hybrid")
.add_attr_option<Array<String>>("keys")
.add_attr_option<Array<String>>("libs")
.add_attr_option<String>("device")
Map<Tensor, Buffer> extern_buffer) {
// Check if current lower target is CUDA
auto target = tvm::Target::Current(true);
- if (target.defined() && target->id->name != "cuda") {
+ if (target.defined() && target->kind->name != "cuda") {
return stmt;
}
if (func->GetAttr<Integer>(tvm::attr::kCallingConv, Integer(CallingConv::kDefault)) ==
CallingConv::kDefault) {
- MemoryAccessVerifier v(func, target.value()->id->device_type);
+ MemoryAccessVerifier v(func, target.value()->kind->device_type);
v.Run();
return !v.Failed();
} else {
auto target = f->GetAttr<Target>(tvm::attr::kTarget);
CHECK(target.defined()) << "LowerCustomDatatypes: Require the target attribute";
- n->body = CustomDatatypesLowerer(target.value()->id->name)(std::move(n->body));
+ n->body = CustomDatatypesLowerer(target.value()->kind->name)(std::move(n->body));
return f;
};
return CreatePrimFuncPass(pass_func, 0, "tir.LowerCustomDatatypes", {});
arith::Analyzer analyzer;
auto mtriple = target.value()->GetAttr<runtime::String>("mtriple", "");
n->body =
- IntrinInjecter(&analyzer, target.value()->id->name, mtriple.value())(std::move(n->body));
+ IntrinInjecter(&analyzer, target.value()->kind->name, mtriple.value())(std::move(n->body));
return f;
};
return CreatePrimFuncPass(pass_func, 0, "tir.LowerIntrin", {});
// Also, the warp/wavefront size differs (64 on rocm, 32 on cuda).
bool is_warp_reduction(const std::vector<DataType>& types) const {
// Only cuda target supports warp reductions.
- if ((target_->id->name != "cuda") && (target_->id->name != "rocm")) return false;
+ if ((target_->kind->name != "cuda") && (target_->kind->name != "rocm")) return false;
// rocm only supports 32 bit operands for shuffling at the moment
- if ((target_->id->name == "rocm") && (std::any_of(types.begin(), types.end(), [](DataType ty) {
+ if ((target_->kind->name == "rocm") &&
+ (std::any_of(types.begin(), types.end(), [](DataType ty) {
if (ty.is_vector()) return true;
return ty.bits() != 32;
}))) {
auto target = func->GetAttr<Target>(tvm::attr::kTarget);
CHECK(target.defined()) << "MakePackedAPI: Require the target attribute";
- int target_device_type = target.value()->id->device_type;
+ int target_device_type = target.value()->kind->device_type;
std::string name_hint = global_symbol.value();
#include <dmlc/logging.h>
#include <gtest/gtest.h>
-#include <tvm/target/target_id.h>
+#include <tvm/target/target_kind.h>
#include <cmath>
#include <string>
using namespace tvm;
-TVM_REGISTER_TARGET_ID("TestTargetId")
+TVM_REGISTER_TARGET_KIND("TestTargetKind")
.set_attr<std::string>("Attr1", "Value1")
.add_attr_option<Bool>("my_bool")
.add_attr_option<Array<String>>("your_names")
.add_attr_option<Map<String, Integer>>("her_maps");
-TEST(TargetId, GetAttrMap) {
- auto map = tvm::TargetId::GetAttrMap<std::string>("Attr1");
- auto target_id = tvm::TargetId::Get("TestTargetId");
- std::string result = map[target_id];
+TEST(TargetKind, GetAttrMap) {
+ auto map = tvm::TargetKind::GetAttrMap<std::string>("Attr1");
+ auto target_kind = tvm::TargetKind::Get("TestTargetKind");
+ std::string result = map[target_kind];
CHECK_EQ(result, "Value1");
}
-TEST(TargetId, SchemaValidation) {
+TEST(TargetKind, SchemaValidation) {
tvm::Map<String, ObjectRef> target;
{
tvm::Array<String> your_names{"junru", "jian"};
target.Set("my_bool", Bool(true));
target.Set("your_names", your_names);
target.Set("her_maps", her_maps);
- target.Set("id", String("TestTargetId"));
+ target.Set("kind", String("TestTargetKind"));
}
TargetValidateSchema(target);
tvm::Map<String, ObjectRef> target_host(target.begin(), target.end());
TargetValidateSchema(target);
}
-TEST(TargetId, SchemaValidationFail) {
+TEST(TargetKind, SchemaValidationFail) {
tvm::Map<String, ObjectRef> target;
{
tvm::Array<String> your_names{"junru", "jian"};
target.Set("your_names", your_names);
target.Set("her_maps", her_maps);
target.Set("ok", ObjectRef(nullptr));
- target.Set("id", String("TestTargetId"));
+ target.Set("kind", String("TestTargetKind"));
}
bool failed = false;
try {
from tvm import te
from tvm.target import cuda, rocm, mali, intel_graphics, arm_cpu, vta, bifrost, hexagon
+
@tvm.target.generic_func
def mygeneric(data):
# default generic function
return data + 1
+
@mygeneric.register(["cuda", "gpu"])
def cuda_func(data):
return data + 2
+
@mygeneric.register("rocm")
def rocm_func(data):
return data + 3
+
@mygeneric.register("cpu")
def rocm_func(data):
return data + 10
def test_target_string_parse():
target = tvm.target.create("cuda -model=unknown -libs=cublas,cudnn")
- assert target.id.name == "cuda"
+ assert target.kind.name == "cuda"
assert target.model == "unknown"
assert set(target.keys) == set(['cuda', 'gpu'])
assert set(target.libs) == set(['cublas', 'cudnn'])
def test_target_create():
- targets = [cuda(), rocm(), mali(), intel_graphics(), arm_cpu('rk3399'), vta(), bifrost()]
+ targets = [cuda(), rocm(), mali(), intel_graphics(),
+ arm_cpu('rk3399'), vta(), bifrost()]
for tgt in targets:
assert tgt is not None
if __name__ == "__main__":
test_target_dispatch()
test_target_string_parse()
- test_target_create()
\ No newline at end of file
+ test_target_create()
projects / platform / upstream / tvm.git / commitdiff