This is an automated email from the ASF dual-hosted git repository.
tqchen pushed a commit to branch main
in repository https://gitbox.apache.org/repos/asf/tvm.git
The following commit(s) were added to refs/heads/main by this push:
new 302aaf9f96 [IR] Rename Var name_hint field to name (#20016)
302aaf9f96 is described below
commit 302aaf9f961a6e0a2c5cc23dc87e51fbaabd1e42
Author: Tianqi Chen <[email protected]>
AuthorDate: Fri Jul 17 05:34:31 2026 +0800
[IR] Rename Var name_hint field to name (#20016)
Rename the reflected local `Var` field from `name_hint` to `name` and
update its typed C++ consumers. Preserve distinct named-node APIs and
the Python constructor keyword compatibility path, while making `.name`
the sole stored Var property. Upgrade legacy compact JSON records for
current and pre-unification Var schemas.
Validation: full runtime/compiler build, focused C++ Var copy-helper
test, focused Python IR/Relax/TIRx/script tests, Vulkan codegen syntax
build, touched-file pre-commit checks, and `git diff --check`.
---
include/tvm/ir/expr.h | 9 ++--
include/tvm/relax/expr.h | 2 +-
include/tvm/s_tir/data_layout.h | 10 ++--
include/tvm/tirx/var.h | 9 ++--
include/tvm/topi/transform.h | 2 +-
python/tvm/ir/expr.py | 9 +---
python/tvm/ir/json_compact.py | 10 ++--
python/tvm/relax/expr.py | 2 +-
src/arith/solve_linear_equation.cc | 2 +-
src/arith/unwrap_vector_expr.cc | 2 +-
src/backend/cuda/codegen/codegen_cuda.cc | 16 +++----
src/backend/vulkan/codegen/codegen_spirv.cc | 12 ++---
src/ir/expr.cc | 12 ++---
src/relax/analysis/var2value.cc | 4 +-
src/relax/backend/contrib/codegen_c/codegen_c.h | 2 +-
.../backend/contrib/codegen_json/codegen_json.h | 2 +-
src/relax/backend/contrib/cutlass/codegen.cc | 4 +-
src/relax/backend/vm/codegen_vm.cc | 2 +-
src/relax/backend/vm/codegen_vm_tir.cc | 2 +-
src/relax/backend/vm/vm_shape_lower.cc | 2 +-
src/relax/distributed/transform/lower_distir.cc | 2 +-
.../distributed/transform/propagate_sharding.cc | 7 ++-
src/relax/ir/binding_rewrite.cc | 2 +-
src/relax/ir/block_builder.cc | 2 +-
src/relax/ir/dataflow_matcher.cc | 4 +-
src/relax/ir/expr.cc | 8 ++--
src/relax/ir/expr_functor.cc | 8 ++--
src/relax/script/builder/frame.cc | 2 +-
src/relax/script/builder/ir.cc | 2 +-
src/relax/script/builder/utils.h | 4 +-
src/relax/script/printer/tir.cc | 2 +-
src/relax/script/printer/utils.h | 2 +-
src/relax/training/utils.cc | 2 +-
src/relax/transform/adjust_matmul_order.cc | 2 +-
src/relax/transform/bind_params.cc | 4 +-
src/relax/transform/bind_symbolic_vars.cc | 4 +-
src/relax/transform/bundle_model_params.cc | 6 +--
src/relax/transform/canonicalize_bindings.cc | 2 +-
src/relax/transform/expand_tuple_arguments.cc | 6 +--
src/relax/transform/fuse_ops.cc | 13 +++---
src/relax/transform/fuse_tir.cc | 2 +-
src/relax/transform/gradient.cc | 17 ++++---
src/relax/transform/inline_functions.cc | 2 +-
src/relax/transform/lambda_lift.cc | 4 +-
src/relax/transform/lazy_transform_params.cc | 6 +--
src/relax/transform/lift_transform_params.cc | 4 +-
src/relax/transform/realize_vdevice.cc | 4 +-
src/relax/transform/remove_unused_parameters.cc | 2 +-
src/relax/transform/rewrite_cuda_graph.cc | 6 +--
.../specialize_primfunc_based_on_callsite.cc | 2 +-
src/relax/transform/static_plan_block_memory.cc | 6 +--
src/relax/transform/to_mixed_precision.cc | 4 +-
src/relax/transform/to_non_dataflow.cc | 2 +-
src/relax/transform/utils.h | 4 +-
src/s_tir/analysis/verify_gpu_code.cc | 4 +-
src/s_tir/data_layout.cc | 26 +++++------
.../schedule/primitive/annotate_buffer_access.cc | 4 +-
src/s_tir/schedule/primitive/cache_index.cc | 2 +-
src/s_tir/schedule/primitive/cache_read_write.cc | 6 +--
.../schedule/primitive/layout_transformation.cc | 8 ++--
.../schedule/primitive/loop_transformation.cc | 22 ++++-----
src/s_tir/schedule/primitive/reduction.cc | 6 +--
src/s_tir/schedule/trace.cc | 2 +-
src/s_tir/schedule/transform.cc | 5 +-
src/s_tir/transform/compact_buffer_region.cc | 2 +-
src/s_tir/transform/convert_blocks_to_opaque.cc | 2 +-
src/s_tir/transform/inject_double_buffer.cc | 4 +-
src/s_tir/transform/inject_virtual_thread.cc | 2 +-
.../transform/lower_cross_thread_reduction.cc | 4 +-
src/s_tir/transform/memhammer_coalesce.cc | 2 +-
.../transform/merge_shared_memory_allocations.cc | 2 +-
src/s_tir/transform/renew_defs.cc | 2 +-
src/s_tir/transform/tensorcore_infer_fragment.cc | 4 +-
src/script/ir_builder/ir/ir.cc | 2 +-
src/target/llvm/codegen_cpu.cc | 9 ++--
src/target/llvm/codegen_llvm.cc | 12 ++---
src/target/source/codegen_source_base.cc | 6 +--
src/te/operation/create_primfunc.cc | 6 +--
src/tirx/analysis/var_use_def_analysis.cc | 4 +-
src/tirx/analysis/verify_well_formed.cc | 14 +++---
src/tirx/ir/buffer.cc | 6 +--
src/tirx/ir/index_map.cc | 11 ++---
src/tirx/script/builder/ir.cc | 4 +-
src/tirx/script/printer/buffer.cc | 2 +-
src/tirx/script/printer/expr.cc | 6 +--
src/tirx/script/printer/stmt.cc | 2 +-
src/tirx/script/printer/utils.h | 2 +-
src/tirx/transform/ir_utils.cc | 4 +-
src/tirx/transform/ir_utils.h | 4 +-
src/tirx/transform/lower_intrin.cc | 2 +-
src/tirx/transform/lower_warp_memory.cc | 2 +-
src/tirx/transform/narrow_datatype.cc | 2 +-
src/tirx/transform/split_host_device.cc | 4 +-
src/tirx/transform/storage_rewrite.cc | 31 ++++++-------
src/tirx/transform/tile_primitive_dispatch.cc | 2 +-
src/tirx/transform/tvm_ffi_binder.cc | 16 +++----
src/tirx/transform/unsupported_dtype_legalize.cc | 8 ++--
src/tirx/transform/update_pointer_storage_scope.cc | 2 +-
src/tirx/transform/vectorize_loop.cc | 10 ++--
tests/cpp/expr_test.cc | 6 +--
tests/python/ir/test_node_reflection.py | 54 +++++++++++++++++-----
tests/python/relax/test_expr.py | 16 +++++--
102 files changed, 325 insertions(+), 297 deletions(-)
diff --git a/include/tvm/ir/expr.h b/include/tvm/ir/expr.h
index 0af9a8a103..4d614586d5 100644
--- a/include/tvm/ir/expr.h
+++ b/include/tvm/ir/expr.h
@@ -268,12 +268,12 @@ TVM_DLL PrimExpr operator~(PrimExpr a);
*/
class VarNode : public ExprNode {
public:
- /*! \brief The hint to the variable name. */
- ffi::String name_hint;
+ /*! \brief The variable name. */
+ ffi::String name;
static void RegisterReflection() {
namespace refl = tvm::ffi::reflection;
- refl::ObjectDef<VarNode>().def_ro("name_hint", &VarNode::name_hint,
+ refl::ObjectDef<VarNode>().def_ro("name", &VarNode::name,
refl::AttachFieldFlag::SEqHashIgnore());
}
@@ -285,8 +285,7 @@ class VarNode : public ExprNode {
/*! \brief Managed reference to VarNode. */
class Var : public Expr {
public:
- TVM_DLL explicit Var(ffi::String name_hint, ffi::Optional<Type>
ty_annotation,
- Span span = Span());
+ TVM_DLL explicit Var(ffi::String name, ffi::Optional<Type> ty_annotation,
Span span = Span());
/*! \brief Return a fresh ordinary Var with the same type and a new name. */
TVM_DLL Var CopyWithName(const ffi::String& name) const;
diff --git a/include/tvm/relax/expr.h b/include/tvm/relax/expr.h
index bb1248016a..d534c3948e 100644
--- a/include/tvm/relax/expr.h
+++ b/include/tvm/relax/expr.h
@@ -148,7 +148,7 @@ class DataflowVarNode : public VarNode {
class DataflowVar : public Var {
public:
- TVM_DLL explicit DataflowVar(ffi::String name_hint, ffi::Optional<Type>
ty_annotation,
+ TVM_DLL explicit DataflowVar(ffi::String name, ffi::Optional<Type>
ty_annotation,
Span span = Span());
TVM_FFI_DEFINE_OBJECT_REF_METHODS_NULLABLE(DataflowVar, Var,
DataflowVarNode);
diff --git a/include/tvm/s_tir/data_layout.h b/include/tvm/s_tir/data_layout.h
index ee6d51832d..4ba7ff5b29 100644
--- a/include/tvm/s_tir/data_layout.h
+++ b/include/tvm/s_tir/data_layout.h
@@ -46,7 +46,7 @@ class SLayoutAxis {
public:
static const SLayoutAxis& Get(const char name);
- // Get the singleton SLayoutAxis using itvar->var->name_hint
+ // Get the singleton SLayoutAxis using itvar->var->name
static const SLayoutAxis& Get(const tirx::IterVar& itvar);
// Get the singleton SLayoutAxis using name[0] (size of name must be 1).
@@ -230,7 +230,7 @@ class SLayout : public ffi::ObjectRef {
for (auto dst_axis : iter_vars) {
if (SLayoutAxis::Get(dst_axis).IsPrimal()) {
if (!this->Contains(SLayoutAxis::Get(dst_axis))) {
- new_src_layout_str += dst_axis->var->name_hint;
+ new_src_layout_str += dst_axis->var->name;
}
}
}
@@ -252,7 +252,7 @@ class SLayout : public ffi::ObjectRef {
if (!this->defined()) return -1;
const auto axes = operator->()->axes;
for (size_t i = 0; i < axes.size(); ++i) {
- if (axes[i]->var->name_hint == axis) return static_cast<int32_t>(i);
+ if (axes[i]->var->name == axis) return static_cast<int32_t>(i);
}
return -1;
}
@@ -273,7 +273,7 @@ class SLayout : public ffi::ObjectRef {
* \param iter the input iter var.
* \return the index or -1 if not found.
*/
- inline int32_t IndexOf(const tirx::IterVar& iter) const { return
IndexOf(iter->var->name_hint); }
+ inline int32_t IndexOf(const tirx::IterVar& iter) const { return
IndexOf(iter->var->name); }
/*!
* \brief Get the factor size of the subordinate axis.
@@ -294,7 +294,7 @@ class SLayout : public ffi::ObjectRef {
for (const tirx::IterVar packed_var : operator->()->axes) {
auto iter_vars = UnpackIterVar(packed_var);
for (auto var : iter_vars) {
- if (var->var->name_hint == axis.name()) {
+ if (var->var->name == axis.name()) {
return true;
}
}
diff --git a/include/tvm/tirx/var.h b/include/tvm/tirx/var.h
index 467b2f7bee..28e322b3e3 100644
--- a/include/tvm/tirx/var.h
+++ b/include/tvm/tirx/var.h
@@ -46,13 +46,12 @@ using Var = tvm::Var;
class PrimVar : public PrimExpr {
public:
/*! \brief Construct a scalar variable directly from a primitive type. */
- explicit PrimVar(ffi::String name_hint, PrimType dtype = PrimType::Int(32),
Span span = Span())
- : PrimExpr(
- Var(std::move(name_hint), std::move(dtype),
std::move(span)).as_or_throw<PrimExpr>()) {}
+ explicit PrimVar(ffi::String name, PrimType dtype = PrimType::Int(32), Span
span = Span())
+ : PrimExpr(Var(std::move(name), std::move(dtype),
std::move(span)).as_or_throw<PrimExpr>()) {}
/*! \brief Construct a scalar variable directly from a checked type
annotation. */
- explicit PrimVar(ffi::String name_hint, Type type_annotation, Span span =
Span())
- : PrimExpr(Var(std::move(name_hint), std::move(type_annotation),
std::move(span))
+ explicit PrimVar(ffi::String name, Type type_annotation, Span span = Span())
+ : PrimExpr(Var(std::move(name), std::move(type_annotation),
std::move(span))
.as_or_throw<PrimExpr>()) {}
/*! \brief Safe widening to a general Var view over the same node. */
diff --git a/include/tvm/topi/transform.h b/include/tvm/topi/transform.h
index 5d56a0ca64..e886c1839a 100644
--- a/include/tvm/topi/transform.h
+++ b/include/tvm/topi/transform.h
@@ -653,7 +653,7 @@ inline PrimExpr DynamicCanonicalizeIndex(PrimExpr index,
PrimExpr extent, PrimEx
auto idx_var = index.as<tvm::tirx::PrimVar>();
auto extent_var = extent.as<tvm::tirx::PrimVar>();
- if (idx_var && extent_var && (*idx_var)->name_hint ==
(*extent_var)->name_hint) {
+ if (idx_var && extent_var && (*idx_var)->name == (*extent_var)->name) {
return index;
}
diff --git a/python/tvm/ir/expr.py b/python/tvm/ir/expr.py
index c1f146da87..4c28b7e7f6 100644
--- a/python/tvm/ir/expr.py
+++ b/python/tvm/ir/expr.py
@@ -374,7 +374,7 @@ class Var(_ExprWithOp):
Parameters
----------
name : str
- The API name of the variable. It is stored internally as a name hint.
+ The name of the variable.
ty : Optional[Type or str]
The exact type of the variable. A string denotes a primitive dtype.
@@ -384,7 +384,7 @@ class Var(_ExprWithOp):
"""
- name_hint: str
+ name: str
span: Span | None
def __init__(
@@ -413,11 +413,6 @@ class Var(_ExprWithOp):
raise TypeError("ty must be a Type or primitive dtype string")
self.__init_handle_by_constructor__(_ffi_api.Var, name, ty, span)
- @property
- def name(self) -> str:
- """The user-facing variable name."""
- return self.name_hint
-
@tvm_ffi.register_object("ir.Range")
class Range(Node, Scriptable):
diff --git a/python/tvm/ir/json_compact.py b/python/tvm/ir/json_compact.py
index 948c064e64..1b65b2488e 100644
--- a/python/tvm/ir/json_compact.py
+++ b/python/tvm/ir/json_compact.py
@@ -86,14 +86,16 @@ def upgrade_json(json_str):
raise ValueError("Legacy json graph format detected, we don't support
it anymore.")
# `ir.Var` is the sole runtime variable node. Keep `tvm.ir.load_json`
- # compatible with the exact pre-unification Relax and TIRx schemas.
- # Rewriting nodes in place preserves node indices and shared references.
+ # compatible with the pre-unification Relax/TIRx schemas and with graphs
+ # written before the canonical Var field was renamed to `name`. Rewriting
+ # nodes in place preserves node indices and shared references.
for node in data.get("nodes", []):
if node.get("type") == "relax.expr.Var":
node["type"] = "ir.Var"
elif node.get("type") == "tirx.Var":
node["type"] = "ir.Var"
+ if node.get("type") in ("ir.Var", "relax.expr.DataflowVar"):
fields = node.get("data", {})
- if "name" in fields and "name_hint" not in fields:
- fields["name_hint"] = fields.pop("name")
+ if "name_hint" in fields and "name" not in fields:
+ fields["name"] = fields.pop("name_hint")
return json.dumps(data, indent=2)
diff --git a/python/tvm/relax/expr.py b/python/tvm/relax/expr.py
index 41b49dd2c5..8c7cdcf992 100644
--- a/python/tvm/relax/expr.py
+++ b/python/tvm/relax/expr.py
@@ -423,7 +423,7 @@ class DataflowVar(Var):
Parameters
----------
name: str
- The API name of the variable. It is stored internally as a name hint.
+ The name of the variable.
ty: Optional[Type]
The type annotation of the variable.
diff --git a/src/arith/solve_linear_equation.cc
b/src/arith/solve_linear_equation.cc
index 382c166305..f2315a473f 100644
--- a/src/arith/solve_linear_equation.cc
+++ b/src/arith/solve_linear_equation.cc
@@ -393,7 +393,7 @@ IntConstraintsTransform SolveLinearEquations(const
IntConstraints& system_to_sol
PrimExpr to_old = analyzer_problem->Simplify(V_inv_x[j]);
std::string name_hint = "n" + std::to_string(new_vars.size());
if (auto old_var = to_old.as<tirx::PrimVar>()) {
- name_hint += "_" + (*old_var)->name_hint;
+ name_hint += "_" + (*old_var)->name;
}
PrimVar v(name_hint, V_inv_x[j].ty());
solution_for_V_inv_x.push_back(v);
diff --git a/src/arith/unwrap_vector_expr.cc b/src/arith/unwrap_vector_expr.cc
index de6956cf27..ddc55d0732 100644
--- a/src/arith/unwrap_vector_expr.cc
+++ b/src/arith/unwrap_vector_expr.cc
@@ -67,7 +67,7 @@ class Scalarizer : public ExprMutator {
TVM_FFI_ICHECK(it == let_var_remap_.end()) << "Duplicate binding of
variable " << op->var;
PrimType var_ty = op->var.as_or_throw<PrimVar>().ty();
- PrimVar new_var(op->var->name_hint + "_scalar", var_ty.WithLanes(1));
+ PrimVar new_var(op->var->name + "_scalar", var_ty.WithLanes(1));
let_var_remap_[op->var.get()] = new_var;
PrimExpr value = this->VisitPrimExpr(op->value);
diff --git a/src/backend/cuda/codegen/codegen_cuda.cc
b/src/backend/cuda/codegen/codegen_cuda.cc
index 5576d8a578..6f951f9789 100644
--- a/src/backend/cuda/codegen/codegen_cuda.cc
+++ b/src/backend/cuda/codegen/codegen_cuda.cc
@@ -198,22 +198,22 @@ class ThreadIdxExtractor : public tirx::StmtVisitor {
void VisitStmt_(const AttrStmtNode* op) final {
if (op->attr_key == tirx::attr::thread_extent) {
IterVar iv = op->node.as_or_throw<IterVar>();
- if (iv->var->name_hint == "threadIdx.x" || iv->thread_tag ==
"threadIdx.x") {
+ if (iv->var->name == "threadIdx.x" || iv->thread_tag == "threadIdx.x") {
threadIdx_x_ext = op->value;
}
- if (iv->var->name_hint == "threadIdx.y" || iv->thread_tag ==
"threadIdx.y") {
+ if (iv->var->name == "threadIdx.y" || iv->thread_tag == "threadIdx.y") {
threadIdx_y_ext = op->value;
}
- if (iv->var->name_hint == "threadIdx.z" || iv->thread_tag ==
"threadIdx.z") {
+ if (iv->var->name == "threadIdx.z" || iv->thread_tag == "threadIdx.z") {
threadIdx_z_ext = op->value;
}
- if (iv->var->name_hint == "clusterCtaIdx.x" || iv->thread_tag ==
"clusterCtaIdx.x") {
+ if (iv->var->name == "clusterCtaIdx.x" || iv->thread_tag ==
"clusterCtaIdx.x") {
clusterCtaIdx_x_ext = op->value;
}
- if (iv->var->name_hint == "clusterCtaIdx.y" || iv->thread_tag ==
"clusterCtaIdx.y") {
+ if (iv->var->name == "clusterCtaIdx.y" || iv->thread_tag ==
"clusterCtaIdx.y") {
clusterCtaIdx_y_ext = op->value;
}
- if (iv->var->name_hint == "clusterCtaIdx.z" || iv->thread_tag ==
"clusterCtaIdx.z") {
+ if (iv->var->name == "clusterCtaIdx.z" || iv->thread_tag ==
"clusterCtaIdx.z") {
clusterCtaIdx_z_ext = op->value;
}
}
@@ -1995,7 +1995,7 @@ void CodeGenCUDA::PrintWmmaScope(const std::string&
scope, const PrimType& t,
std::stringstream type;
PrintType(t, type);
TVM_FFI_ICHECK(fragment_shapes.count(variable))
- << "Cannot find shape of the wmma fragment " << variable->name_hint;
+ << "Cannot find shape of the wmma fragment " << variable->name;
std::string shape_str = fragment_shapes.at(variable);
if ((t.MatchesCode(DLDataTypeCode::kDLInt, DLDataTypeCode::kDLUInt)) &&
t.bits() < 8 &&
t.lanes() == 1) {
@@ -2047,7 +2047,7 @@ int stoi(const std::string& str) {
int32_t CodeGenCUDA::GetWmmaFragmentSize(const std::string& scope, const
VarNode* variable,
int32_t size) {
TVM_FFI_ICHECK(fragment_shapes.count(variable))
- << "Cannot find shape of the wmma fragment " << variable->name_hint;
+ << "Cannot find shape of the wmma fragment " << variable->name;
std::string shape_str = fragment_shapes.at(variable);
std::pair<int32_t, int32_t> dim = GetWmmaFragmentDimSize(shape_str, scope);
if (dim.first * dim.second != 0)
diff --git a/src/backend/vulkan/codegen/codegen_spirv.cc
b/src/backend/vulkan/codegen/codegen_spirv.cc
index e0d44441a3..7dbb23820b 100644
--- a/src/backend/vulkan/codegen/codegen_spirv.cc
+++ b/src/backend/vulkan/codegen/codegen_spirv.cc
@@ -88,8 +88,8 @@ runtime::SPIRVShader CodeGenSPIRV::BuildFunction(const
PrimFunc& f, const std::s
}
spirv::Value arg_value =
builder_->BufferArgument(builder_->GetSType(value_storage_type),
descriptor_set,
i_buffer++);
- builder_->SetName(arg_value, arg->name_hint);
- storage_info_[arg.get()].SetContentType(value_storage_type,
arg->name_hint);
+ builder_->SetName(arg_value, arg->name);
+ storage_info_[arg.get()].SetContentType(value_storage_type, arg->name);
var_map_[arg.get()] = arg_value;
} else {
PrimType pod_type = arg->ty.as_or_throw<PrimType>();
@@ -207,7 +207,7 @@ spirv::Value CodeGenSPIRV::CreateStorageSync(const
CallNode* op) {
spirv::Value CodeGenSPIRV::VisitExpr_(const VarNode* op) {
auto it = var_map_.find(op);
- TVM_FFI_ICHECK(it != var_map_.end()) << "cannot find variable " <<
op->name_hint;
+ TVM_FFI_ICHECK(it != var_map_.end()) << "cannot find variable " << op->name;
return it->second;
}
@@ -628,8 +628,8 @@ spirv::Value CodeGenSPIRV::VisitExpr_(const BufferLoadNode*
op) {
} else {
TVM_FFI_THROW(InternalError) << "Cannot perform buffer access of buffer
variable '"
- << buffer_var->name_hint << "' with element
type "
- << info.element_type << " using index of type
" << prim_index.ty()
+ << buffer_var->name << "' with element type "
<< info.element_type
+ << " using index of type " << prim_index.ty()
<< " to produce output of type " <<
op->ty.as_or_throw<PrimType>();
return spirv::Value();
}
@@ -706,7 +706,7 @@ void CodeGenSPIRV::VisitStmt_(const BufferStoreNode* op) {
} else {
TVM_FFI_THROW(InternalError) << "Cannot store value of type " << value_type
- << " into buffer variable '" <<
buffer_var->name_hint
+ << " into buffer variable '" <<
buffer_var->name
<< "' with element type " << info.element_type
<< " using index of type " << index_type;
}
diff --git a/src/ir/expr.cc b/src/ir/expr.cc
index 2330e643cd..d44cdb06f6 100644
--- a/src/ir/expr.cc
+++ b/src/ir/expr.cc
@@ -251,9 +251,9 @@ TVM_FFI_STATIC_INIT_BLOCK() {
});
}
-Var::Var(ffi::String name_hint, ffi::Optional<Type> ty_annotation, Span span) {
+Var::Var(ffi::String name, ffi::Optional<Type> ty_annotation, Span span) {
ffi::ObjectPtr<VarNode> n = ffi::make_object<VarNode>();
- n->name_hint = std::move(name_hint);
+ n->name = std::move(name);
if (ty_annotation.has_value()) {
n->ty = ty_annotation.value();
}
@@ -265,12 +265,12 @@ Var Var::CopyWithName(const ffi::String& name) const {
TVM_FFI_CHECK_EQ(type_index(), VarNode::RuntimeTypeIndex(), TypeError)
<< "Cannot copy a Var runtime subtype as an ordinary Var";
ffi::ObjectPtr<VarNode> copy = ffi::make_object<VarNode>(*get());
- copy->name_hint = name;
+ copy->name = name;
return Var(std::move(copy));
}
Var Var::CopyWithSuffix(const ffi::String& suffix) const {
- return CopyWithName(get()->name_hint + suffix);
+ return CopyWithName(get()->name + suffix);
}
Var Var::CopyWithDType(PrimType dtype) const {
@@ -305,8 +305,8 @@ Call::Call(Type ret_ty, Expr op, ffi::Array<Expr> args,
Attrs attrs, ffi::Array<
TVM_FFI_STATIC_INIT_BLOCK() {
namespace refl = tvm::ffi::reflection;
refl::GlobalDef()
- .def("ir.Var", [](ffi::String name_hint, ffi::Optional<Type>
ty_annotation,
- Span span) { return Var(name_hint, ty_annotation,
span); })
+ .def("ir.Var", [](ffi::String name, ffi::Optional<Type> ty_annotation,
+ Span span) { return Var(name, ty_annotation, span); })
.def("ir.GlobalVar", [](ffi::String name) { return GlobalVar(name); })
.def("ir.Call",
[](Type ret_ty, Expr op, ffi::Array<Expr> args, Attrs attrs,
ffi::Array<Type> ty_args,
diff --git a/src/relax/analysis/var2value.cc b/src/relax/analysis/var2value.cc
index a5e30c8139..e18d0e7fca 100644
--- a/src/relax/analysis/var2value.cc
+++ b/src/relax/analysis/var2value.cc
@@ -74,12 +74,12 @@ class Name2BindingAnalysis : public relax::ExprVisitor {
// so we use standard container for internal usage.
std::map<ffi::String, ffi::Array<Binding>> name2bindings_;
void VisitBinding_(const VarBindingNode* binding) override {
- const auto& vname = binding->var->name_hint;
+ const auto& vname = binding->var->name;
name2bindings_[vname].push_back(ffi::GetRef<VarBinding>(binding));
}
void VisitBinding_(const MatchCastNode* binding) override {
- const auto& vname = binding->var->name_hint;
+ const auto& vname = binding->var->name;
name2bindings_[vname].push_back(ffi::GetRef<MatchCast>(binding));
}
};
diff --git a/src/relax/backend/contrib/codegen_c/codegen_c.h
b/src/relax/backend/contrib/codegen_c/codegen_c.h
index 58f76c92a8..8b386f3bfa 100644
--- a/src/relax/backend/contrib/codegen_c/codegen_c.h
+++ b/src/relax/backend/contrib/codegen_c/codegen_c.h
@@ -282,7 +282,7 @@ class CodegenCBase {
for (const auto& arg : args) {
const auto& dtype_str = GetDtypeString(arg);
- code_stream_ << dtype_str << "* " << arg->name_hint << ", ";
+ code_stream_ << dtype_str << "* " << arg->name << ", ";
}
for (size_t i = 0; i < outs.size() - 1; ++i) {
code_stream_ << outs[i].dtype << "* out" << i << ", ";
diff --git a/src/relax/backend/contrib/codegen_json/codegen_json.h
b/src/relax/backend/contrib/codegen_json/codegen_json.h
index 7c9f1a1ea7..8a96a788f4 100644
--- a/src/relax/backend/contrib/codegen_json/codegen_json.h
+++ b/src/relax/backend/contrib/codegen_json/codegen_json.h
@@ -238,7 +238,7 @@ class JSONSerializer : public
relax::MemoizedExprTranslator<NodeEntries> {
void serialize(Function func) {
// First we convert all the parameters into input nodes.
for (const auto& param : func->params) {
- auto node_ptr = std::make_shared<JSONGraphNode>(param->name_hint,
"input" /* op_type_ */);
+ auto node_ptr = std::make_shared<JSONGraphNode>(param->name, "input" /*
op_type_ */);
memo_[param] = AddNode(node_ptr, param);
}
heads_ = VisitExpr(func->body);
diff --git a/src/relax/backend/contrib/cutlass/codegen.cc
b/src/relax/backend/contrib/cutlass/codegen.cc
index 4dacc019d8..cb1d83b1a3 100644
--- a/src/relax/backend/contrib/cutlass/codegen.cc
+++ b/src/relax/backend/contrib/cutlass/codegen.cc
@@ -159,7 +159,7 @@ class CodegenCutlass : public
relax::MemoizedExprTranslator<OutputType>,
void AddParm(Var param) {
ext_func_args_.push_back(param);
- auto v_name = name_sup_->FreshName(param->name_hint);
+ auto v_name = name_sup_->FreshName(param->name);
var_name_map_[param.get()] = v_name;
}
@@ -332,7 +332,7 @@ class CodegenCutlass : public
relax::MemoizedExprTranslator<OutputType>,
/*!
* \brief A mapping from a variable to its unique name.
* We use this since sometimes different parameters to the same function end
up having the same
- * name_hint.
+ * name.
*/
std::unordered_map<const VarNode*, std::string> var_name_map_;
/*! \brief A unique name supply to generate a unique name for each
parameter. */
diff --git a/src/relax/backend/vm/codegen_vm.cc
b/src/relax/backend/vm/codegen_vm.cc
index 176d28eabe..a7dd9a21a8 100644
--- a/src/relax/backend/vm/codegen_vm.cc
+++ b/src/relax/backend/vm/codegen_vm.cc
@@ -91,7 +91,7 @@ class CodeGenVM : public ExprFunctor<Instruction::Arg(const
Expr&)> {
ffi::Array<ffi::String> param_names;
for (Var param : func->params) {
- param_names.push_back(param->name_hint);
+ param_names.push_back(param->name);
}
builder_->EmitFunction(gsymbol.value(), func->params.size(), param_names);
diff --git a/src/relax/backend/vm/codegen_vm_tir.cc
b/src/relax/backend/vm/codegen_vm_tir.cc
index 81482b0293..48191f32a3 100644
--- a/src/relax/backend/vm/codegen_vm_tir.cc
+++ b/src/relax/backend/vm/codegen_vm_tir.cc
@@ -171,7 +171,7 @@ class CodeGenVMTIR : public
ExprFunctor<ffi::Optional<Expr>(const Expr&)> {
ffi::Array<ffi::String> param_names;
for (Var param : func->params) {
- param_names.push_back(param->name_hint);
+ param_names.push_back(param->name);
}
// declare this function.
builder_->DeclareFunction(gsymbol.value(),
vm::VMFuncInfo::FuncKind::kVMTIRFunc);
diff --git a/src/relax/backend/vm/vm_shape_lower.cc
b/src/relax/backend/vm/vm_shape_lower.cc
index 12ff275784..b83a5ed72d 100644
--- a/src/relax/backend/vm/vm_shape_lower.cc
+++ b/src/relax/backend/vm/vm_shape_lower.cc
@@ -368,7 +368,7 @@ class VMShapeLowerMutator
Type ty = GetType(func->params[i]);
std::ostringstream err_ctx;
err_ctx << "ErrorContext(fn=" << gvar->name_hint << ", loc=param[" << i
- << "], param=" << func->params[i]->name_hint << ",
annotation=" << ty << ") ";
+ << "], param=" << func->params[i]->name << ", annotation=" <<
ty << ") ";
this->CheckMatchCast(ty, func->params[i], true, i >= num_input,
err_ctx.str(),
&match_todos);
if (PrimExprSlot* slot = GetPrimValueSlot(func->params[i])) {
diff --git a/src/relax/distributed/transform/lower_distir.cc
b/src/relax/distributed/transform/lower_distir.cc
index 01522183bb..ae0c30c3d5 100644
--- a/src/relax/distributed/transform/lower_distir.cc
+++ b/src/relax/distributed/transform/lower_distir.cc
@@ -116,7 +116,7 @@ class DistIRSharder : public ExprMutator {
Type old_ty = GetType(input);
Type new_ty = ConvertType(old_ty, false);
if (const auto* var = input.as<VarNode>()) {
- Var new_param(var->name_hint, new_ty);
+ Var new_param(var->name, new_ty);
return new_param;
} else if (const auto* constant = input.as<ConstantNode>()) {
for (const auto& spec :
old_ty.as_or_throw<DTensorType>()->placement->dim_specs) {
diff --git a/src/relax/distributed/transform/propagate_sharding.cc
b/src/relax/distributed/transform/propagate_sharding.cc
index 738d5dc0b2..2abeccaaef 100644
--- a/src/relax/distributed/transform/propagate_sharding.cc
+++ b/src/relax/distributed/transform/propagate_sharding.cc
@@ -285,8 +285,7 @@ class ShardingConflictHandler : public ExprVisitor {
if (device_mesh.has_value()) {
TVM_FFI_ICHECK(ffi::StructuralEqual()(device_mesh.value(),
sharding_spec.first))
- << "Sharding conflict detected for tensor " << var->name_hint
- << ": Device Mesh mismatch"
+ << "Sharding conflict detected for tensor " << var->name << ":
Device Mesh mismatch"
<< ". Conflict Handling logic will be added in the future.";
} else {
device_mesh = sharding_spec.first;
@@ -294,7 +293,7 @@ class ShardingConflictHandler : public ExprVisitor {
if (i >= 0) {
int sharding_dim = sharding_spec.second;
TVM_FFI_ICHECK(sharded_mesh_dim.count(sharding_dim) == 0)
- << "Sharding conflict detected for tensor " << var->name_hint
+ << "Sharding conflict detected for tensor " << var->name
<< ": Replicate sharding device mesh axis " << sharding_dim
<< ". Conflict Handling logic will be added in the future.";
sharded_mesh_dim.insert(sharding_dim);
@@ -401,7 +400,7 @@ class DistributedIRBuilder : public ExprMutator {
}
if (const auto* var = tensor.as<VarNode>()) {
- Var new_param(var->name_hint, new_ty);
+ Var new_param(var->name, new_ty);
return new_param;
} else if (const auto* constant = tensor.as<ConstantNode>()) {
Constant new_constant(constant->data, new_ty);
diff --git a/src/relax/ir/binding_rewrite.cc b/src/relax/ir/binding_rewrite.cc
index 44348e6cef..359e1b1652 100644
--- a/src/relax/ir/binding_rewrite.cc
+++ b/src/relax/ir/binding_rewrite.cc
@@ -55,7 +55,7 @@ DataflowBlockRewrite::DataflowBlockRewrite(DataflowBlock dfb,
Function root_fn)
n->to_users_ = std::move(p.first);
n->fn_outputs_ = std::move(p.second);
n->name_supply_ = UniqueNameSupply(n->to_users_.begin(), n->to_users_.end(),
- [](const auto& p) { return
p.first->name_hint; });
+ [](const auto& p) { return p.first->name;
});
data_ = std::move(n);
}
diff --git a/src/relax/ir/block_builder.cc b/src/relax/ir/block_builder.cc
index e1a64d191e..ec722df110 100644
--- a/src/relax/ir/block_builder.cc
+++ b/src/relax/ir/block_builder.cc
@@ -589,7 +589,7 @@ class Normalizer : public BlockBuilderImpl, private
ExprFunctor<Expr(const Expr&
Expr VisitVar_(const typename T::ContainerType* var) {
// Parameters and free-vars must be present with type
// Other vars must have already been normalized through binding
- TVM_FFI_ICHECK(!var->ty.IsMissing()) << "Var " << var->name_hint << " does
not have type.";
+ TVM_FFI_ICHECK(!var->ty.IsMissing()) << "Var " << var->name << " does not
have type.";
return ffi::GetRef<Var>(var);
}
diff --git a/src/relax/ir/dataflow_matcher.cc b/src/relax/ir/dataflow_matcher.cc
index 21717bdb8c..5a982c500e 100644
--- a/src/relax/ir/dataflow_matcher.cc
+++ b/src/relax/ir/dataflow_matcher.cc
@@ -463,7 +463,7 @@ PrimExpr DFPatternMatcher::SimplifyCondition(PrimExpr
condition) {
auto sort_key = [](PrimExpr expr) -> ffi::String {
if (const auto* equal = expr.as<tirx::EQNode>()) {
if (auto var = equal->a.as<tirx::PrimVar>()) {
- return var.value()->name_hint;
+ return var.value()->name;
}
}
return "";
@@ -582,7 +582,7 @@ bool DFPatternMatcher::VisitDFPattern_(const
VarPatternNode* op, const Expr& exp
// We don't jump for var pattern, as there's no need to access its value to
judge it.
if (const auto* var_node = expr.as<VarNode>()) {
// "" means any name.
- return "" == op->name_hint() || op->name_hint() == var_node->name_hint;
+ return "" == op->name_hint() || op->name_hint() == var_node->name;
}
return false;
}
diff --git a/src/relax/ir/expr.cc b/src/relax/ir/expr.cc
index 050b2dc8b0..3242fb2224 100644
--- a/src/relax/ir/expr.cc
+++ b/src/relax/ir/expr.cc
@@ -138,9 +138,9 @@ TVM_FFI_STATIC_INIT_BLOCK() {
});
}
-DataflowVar::DataflowVar(ffi::String name_hint, ffi::Optional<Type>
ty_annotation, Span span) {
+DataflowVar::DataflowVar(ffi::String name, ffi::Optional<Type> ty_annotation,
Span span) {
ffi::ObjectPtr<DataflowVarNode> n = ffi::make_object<DataflowVarNode>();
- n->name_hint = std::move(name_hint);
+ n->name = std::move(name);
if (ty_annotation.has_value()) {
n->ty = ty_annotation.value();
}
@@ -151,8 +151,8 @@ DataflowVar::DataflowVar(ffi::String name_hint,
ffi::Optional<Type> ty_annotatio
TVM_FFI_STATIC_INIT_BLOCK() {
namespace refl = tvm::ffi::reflection;
refl::GlobalDef().def("relax.DataflowVar",
- [](ffi::String name_hint, ffi::Optional<Type>
ty_annotation, Span span) {
- return DataflowVar(name_hint, ty_annotation, span);
+ [](ffi::String name, ffi::Optional<Type>
ty_annotation, Span span) {
+ return DataflowVar(name, ty_annotation, span);
});
}
diff --git a/src/relax/ir/expr_functor.cc b/src/relax/ir/expr_functor.cc
index 4647d51070..3746acd758 100644
--- a/src/relax/ir/expr_functor.cc
+++ b/src/relax/ir/expr_functor.cc
@@ -928,7 +928,7 @@ Var ExprMutator::VisitVarDef_(const DataflowVarNode* var) {
// where we should produce a DataflowVar.
if (!output->IsInstance<DataflowVarNode>()) {
Var delegated_output = output;
- output = DataflowVar(output->name_hint, GetType(output), output->span);
+ output = DataflowVar(output->name, GetType(output), output->span);
var_remap_[delegated_output] = output;
}
return output;
@@ -940,7 +940,7 @@ Var ExprMutator::VisitVarDef_(const VarNode* var) {
if (ty.same_as(var->ty)) {
return ffi::GetRef<Var>(var);
} else {
- return Var(var->name_hint, ty, var->span);
+ return Var(var->name, ty, var->span);
}
} else {
return ffi::GetRef<Var>(var);
@@ -1041,8 +1041,8 @@ Var ExprMutator::WithType(Var var, Type ty) {
if (var->ty.same_as(ty) || ffi::StructuralEqual()(var->ty, ty)) {
return var;
} else {
- Var new_var = var.as<DataflowVarNode>() ? DataflowVar(var->name_hint,
ty, var->span)
- : Var(var->name_hint, ty,
var->span);
+ Var new_var = var.as<DataflowVarNode>() ? DataflowVar(var->name, ty,
var->span)
+ : Var(var->name, ty, var->span);
return new_var;
}
} else {
diff --git a/src/relax/script/builder/frame.cc
b/src/relax/script/builder/frame.cc
index edee981cc6..798ce96f7d 100644
--- a/src/relax/script/builder/frame.cc
+++ b/src/relax/script/builder/frame.cc
@@ -169,7 +169,7 @@ void BindingBlockFrameNode::ExitWithScope() {
ffi::Array<tvm::Var> new_output_vars;
std::unordered_map<tvm::Var, tvm::Var, ffi::ObjectPtrHash,
ffi::ObjectPtrEqual> var_remap;
for (const auto& output_var : output_vars) {
- tvm::Var new_output_var(output_var->name_hint,
tvm::relax::GetType(output_var));
+ tvm::Var new_output_var(output_var->name,
tvm::relax::GetType(output_var));
new_output_vars.push_back(new_output_var);
var_remap[output_var] = new_output_var;
}
diff --git a/src/relax/script/builder/ir.cc b/src/relax/script/builder/ir.cc
index 87115eca5f..ba047c6e3f 100644
--- a/src/relax/script/builder/ir.cc
+++ b/src/relax/script/builder/ir.cc
@@ -38,7 +38,7 @@ TVM_STATIC_IR_FUNCTOR(Namer, vtable)
ffi::String name) -> void {
using tvm::relax::DataflowVarNode;
DataflowVarNode* var =
const_cast<DataflowVarNode*>(node.as<DataflowVarNode>());
- var->name_hint = name;
+ var->name = name;
});
/////////////////////////////// Function ////////////////////////////////
diff --git a/src/relax/script/builder/utils.h b/src/relax/script/builder/utils.h
index bc8f087a0d..e4d63c62f6 100644
--- a/src/relax/script/builder/utils.h
+++ b/src/relax/script/builder/utils.h
@@ -99,7 +99,7 @@ inline tvm::relax::SeqExpr GetSeqExprForBranch(const
SeqExprFrame& frame, ffi::S
TVM_FFI_ICHECK(!last_binding->var->IsInstance<tvm::relax::DataflowVarNode>())
<< "A non-dataflow var is expected in the last binding of '" << method
<< "'.";
- *var_name = last_binding->var->name_hint;
+ *var_name = last_binding->var->name;
// Step 3. Re-collect binding blocks to replace the last binding.
ffi::Array<tvm::relax::BindingBlock>
new_blocks(frame->binding_blocks.begin(),
@@ -107,7 +107,7 @@ inline tvm::relax::SeqExpr GetSeqExprForBranch(const
SeqExprFrame& frame, ffi::S
ffi::Array<tvm::relax::Binding>
last_block_bindings(last_block->bindings.begin(),
last_block->bindings.end() - 1);
- tvm::Var new_var(last_binding->var->name_hint + output_var_suffix,
+ tvm::Var new_var(last_binding->var->name + output_var_suffix,
tvm::relax::GetType(last_binding->var));
tvm::relax::Expr body;
diff --git a/src/relax/script/printer/tir.cc b/src/relax/script/printer/tir.cc
index 4b989958a5..dfd742a0b1 100644
--- a/src/relax/script/printer/tir.cc
+++ b/src/relax/script/printer/tir.cc
@@ -62,7 +62,7 @@ Doc PrintCanonicalVar(Var n, AccessPath n_p, IRDocsifier d) {
TVM_FFI_ICHECK(f->is_func);
f->func_vars->insert(n.get());
}
- IdDoc var = d->Define(n, ffi::GetRef<Frame>(f), n->name_hint.empty() ? "v"
: n->name_hint);
+ IdDoc var = d->Define(n, ffi::GetRef<Frame>(f), n->name.empty() ? "v" :
n->name);
var->source_paths.push_back(n_p);
f->stmts.push_back(AssignDoc(var, PrintVarCreation(prim_var, n_p, d),
std::nullopt));
}
diff --git a/src/relax/script/printer/utils.h b/src/relax/script/printer/utils.h
index 646f8157f3..6fe09a672b 100644
--- a/src/relax/script/printer/utils.h
+++ b/src/relax/script/printer/utils.h
@@ -75,7 +75,7 @@ inline std::string ReprPrintRelax(const ffi::ObjectRef& obj,
const PrinterConfig
}
inline IdDoc DefineRelaxVar(const tvm::Var& var, const Frame& frame, const
IRDocsifier& d) {
- return d->Define(var, frame, var->name_hint.empty() ? "v" : var->name_hint);
+ return d->Define(var, frame, var->name.empty() ? "v" : var->name);
}
inline ffi::Optional<ExprDoc> TypeAsAnn(const tvm::Var& v, const AccessPath&
v_p,
diff --git a/src/relax/training/utils.cc b/src/relax/training/utils.cc
index aca5ed16a9..39365e84ae 100644
--- a/src/relax/training/utils.cc
+++ b/src/relax/training/utils.cc
@@ -192,7 +192,7 @@ class AppendLossMutator : private ExprMutator {
for (int i = 0; i < num_backbone_outputs_; ++i) {
auto var = backbone_return_arr_[i];
if (other_outputs_var.count(var) == 0 && !var->ty.as<PrimTypeNode>()) {
- auto new_var = DataflowVar(var->name_hint, GetType(var), var->span);
+ auto new_var = DataflowVar(var->name, GetType(var), var->span);
this->var_remap_[var] = new_var;
backbone_return_arr_.Set(i, new_var);
}
diff --git a/src/relax/transform/adjust_matmul_order.cc
b/src/relax/transform/adjust_matmul_order.cc
index e93063ae28..1f0cf87281 100644
--- a/src/relax/transform/adjust_matmul_order.cc
+++ b/src/relax/transform/adjust_matmul_order.cc
@@ -102,7 +102,7 @@ std::tuple<DFPattern, ffi::TypedFunction<Expr(Expr,
ffi::Map<DFPattern, Expr>)>>
if (upper_bounds || lower_bounds) {
ffi::Map<ffi::String, tirx::Var> name_lookup;
for (const auto& tir_var : TIRVarsInType(GetType(func))) {
- name_lookup.Set(tir_var->name_hint, tir_var);
+ name_lookup.Set(tir_var->name, tir_var);
symbolic_var_constraints = symbolic_var_constraints && (0 <=
tir_var.as_or_throw<PrimExpr>());
}
diff --git a/src/relax/transform/bind_params.cc
b/src/relax/transform/bind_params.cc
index 1866584041..bbe15e7b5e 100644
--- a/src/relax/transform/bind_params.cc
+++ b/src/relax/transform/bind_params.cc
@@ -39,7 +39,7 @@ ffi::Map<Var, Expr> NormalizeBindings(const Function& func,
std::unordered_map<std::string, ffi::Array<tvm::Var>> string_lookup;
std::unordered_set<const tvm::VarNode*> var_set;
for (const auto& param : func->params) {
- string_lookup[param->name_hint].push_back(param);
+ string_lookup[param->name].push_back(param);
var_set.insert(param.get());
}
@@ -52,7 +52,7 @@ ffi::Map<Var, Expr> NormalizeBindings(const Function& func,
TVM_FFI_ICHECK(it != string_lookup.end())
<< "Function does not have parameter with name \"" << str << "\". "
<< "Function parameters are named "
- << func->params.Map([](const auto& param) { return param->name_hint;
});
+ << func->params.Map([](const auto& param) { return param->name; });
TVM_FFI_ICHECK_EQ(it->second.size(), 1)
<< "Function contains multiple parameters with name \"" << str <<
"\". "
<< "The Relax variables " << it->second << " are all named \"" <<
str << "\"";
diff --git a/src/relax/transform/bind_symbolic_vars.cc
b/src/relax/transform/bind_symbolic_vars.cc
index b0ed4a3674..05a898d7c7 100644
--- a/src/relax/transform/bind_symbolic_vars.cc
+++ b/src/relax/transform/bind_symbolic_vars.cc
@@ -43,7 +43,7 @@ Function FunctionBindSymbolicVars(
std::unordered_map<std::string, ffi::Array<tirx::Var>> string_lookup;
std::unordered_set<const tirx::VarNode*> symbolic_var_set;
for (const auto& var : old_symbolic_vars) {
- string_lookup[var->name_hint].push_back(var);
+ string_lookup[var->name].push_back(var);
symbolic_var_set.insert(var.get());
}
@@ -107,7 +107,7 @@ IRModule ModuleBindSymbolicVars(
std::unordered_set<std::string> var_names;
std::unordered_set<const tirx::VarNode*> vars;
for (const auto& var : DefinedSymbolicVars(func)) {
- var_names.insert(var->name_hint);
+ var_names.insert(var->name);
vars.insert(var.get());
}
diff --git a/src/relax/transform/bundle_model_params.cc
b/src/relax/transform/bundle_model_params.cc
index c2de36d9d3..417400fb1a 100644
--- a/src/relax/transform/bundle_model_params.cc
+++ b/src/relax/transform/bundle_model_params.cc
@@ -119,7 +119,7 @@ class ModelParamBundler : public ExprMutator {
for (const Var& var : prim_params) {
auto it = var_to_expr_.find(var);
TVM_FFI_ICHECK(it != var_to_expr_.end());
- var_remap_[var] = builder_->Emit((*it).second, var->name_hint);
+ var_remap_[var] = builder_->Emit((*it).second, var->name);
}
BindingBlock prologue = builder_->EndBlock();
@@ -150,8 +150,8 @@ class ModelParamBundler : public ExprMutator {
TVM_FFI_ICHECK(field_type != var_to_field_type_.end());
Type rebound_type = VisitExprDepTypeField(GetType(var));
Var replacement = (*field_type).second.same_as(rebound_type)
- ? builder_->Emit((*it).second, op->name_hint)
- : builder_->EmitMatchCast((*it).second,
rebound_type, op->name_hint);
+ ? builder_->Emit((*it).second, op->name)
+ : builder_->EmitMatchCast((*it).second,
rebound_type, op->name);
return replacement;
}
return ExprMutator::VisitExpr_(op);
diff --git a/src/relax/transform/canonicalize_bindings.cc
b/src/relax/transform/canonicalize_bindings.cc
index 65851b0cb3..7d1d168fe7 100644
--- a/src/relax/transform/canonicalize_bindings.cc
+++ b/src/relax/transform/canonicalize_bindings.cc
@@ -233,7 +233,7 @@ class CanonicalizePlanner : public ExprVisitor {
// of trivial bindings, then we can replace it with a DataflowVar.
for (auto var : visitor.defined_inside_dataflow_) {
if (!var.as<DataflowVarNode>() &&
!visitor.used_outside_home_dataflow_.count(var)) {
- DataflowVar new_var(var->name_hint, GetType(var));
+ DataflowVar new_var(var->name, GetType(var));
plan.replace_binding.Set(var, new_var);
plan.replace_usage.Set(var, new_var);
diff --git a/src/relax/transform/expand_tuple_arguments.cc
b/src/relax/transform/expand_tuple_arguments.cc
index fa37f9ca07..8d3603bcc2 100644
--- a/src/relax/transform/expand_tuple_arguments.cc
+++ b/src/relax/transform/expand_tuple_arguments.cc
@@ -49,9 +49,9 @@ ffi::Optional<Function> ExpandParams(Function func) {
if (auto ty = param->ty.as<TupleTypeNode>()) {
ffi::Array<Expr> internal_tuple;
for (size_t i = 0; i < ty->fields.size(); i++) {
- auto name = static_cast<const std::stringstream&>(std::stringstream()
- << param->name_hint
<< "_" << i)
- .str();
+ auto name =
+ static_cast<const std::stringstream&>(std::stringstream() <<
param->name << "_" << i)
+ .str();
Var new_param(name, ty->fields[i]);
internal_tuple.push_back(new_param);
expand_param(new_param);
diff --git a/src/relax/transform/fuse_ops.cc b/src/relax/transform/fuse_ops.cc
index 7d8986218b..d2f9870a38 100644
--- a/src/relax/transform/fuse_ops.cc
+++ b/src/relax/transform/fuse_ops.cc
@@ -503,7 +503,7 @@ class FunctionCreator : public ExprMutator {
TVM_FFI_ICHECK(!item_indices.empty());
int param_idx = tuple_param_idx_[tuple_arg];
Var param = params_[param_idx];
- ffi::String param_name = params_[param_idx]->name_hint;
+ ffi::String param_name = params_[param_idx]->name;
TupleType param_ty = tuple_arg->ty.as_or_throw<TupleType>();
ffi::Array<Expr> item_args;
@@ -639,9 +639,8 @@ class FunctionCreator : public ExprMutator {
}
if ((var == nullptr || defined_vars_.count(var) == 0) &&
(lift_constant_ || !expr->IsInstance<ConstantNode>())) {
- ffi::String name = var != nullptr
- ? var->name_hint
- : ffi::String("param_" +
std::to_string(n_param_for_const_++));
+ ffi::String name =
+ var != nullptr ? var->name : ffi::String("param_" +
std::to_string(n_param_for_const_++));
Type param_ty = GetType(expr);
if (!IsInlinableConstants(expr)) {
Var param(std::move(name), GetType(expr));
@@ -956,8 +955,8 @@ class OperatorFusor : public ExprMutator {
if (producer_group != cur_group) {
for (Group* depgroup : group_deps_[producer_group]) {
TVM_FFI_ICHECK(depgroup != cur_group)
- << "A cyclic dependency detected between the groups " <<
binding->var->name_hint
- << " and " << used_var->name_hint << " are in.";
+ << "A cyclic dependency detected between the groups " <<
binding->var->name
+ << " and " << used_var->name << " are in.";
}
group_deps_[cur_group].push_back(producer_group);
}
@@ -1355,7 +1354,7 @@ class CompositeFunctionAnnotator : public ExprMutator {
ffi::Array<Expr> params;
for (auto v : func_node->params) {
- Var new_v(v->name_hint, GetType(v));
+ Var new_v(v->name, GetType(v));
param_vars.push_back(new_v);
params.push_back(new_v);
}
diff --git a/src/relax/transform/fuse_tir.cc b/src/relax/transform/fuse_tir.cc
index e74584856d..3a653ea753 100644
--- a/src/relax/transform/fuse_tir.cc
+++ b/src/relax/transform/fuse_tir.cc
@@ -931,7 +931,7 @@ class FusedTIRConstructor : public ExprVisitor {
<< "All tuple parameters should be expanded before this point in
FuseTIR. "
<< "However, parameter " << relax_param << " has type " << ty;
- auto name_hint = relax_param->name_hint;
+ auto name_hint = relax_param->name;
if (const auto* tensor = ty.as<TensorTypeNode>()) {
// Case 1. The relax param is a Tensor, we directly create a tirx var
and buffer
diff --git a/src/relax/transform/gradient.cc b/src/relax/transform/gradient.cc
index f11d6af8c2..a927014ea4 100644
--- a/src/relax/transform/gradient.cc
+++ b/src/relax/transform/gradient.cc
@@ -166,7 +166,7 @@ class CheckpointCollector : private ExprMutator {
// Add remapping from binding->var to new_var
if (!binding->var.as<DataflowVarNode>() &&
var->IsInstance<DataflowVarNode>()) {
// For output binding, emit a dummy binding
- this->var_remap_[binding->var] = builder_->EmitOutput(orig_var,
orig_var->name_hint);
+ this->var_remap_[binding->var] = builder_->EmitOutput(orig_var,
orig_var->name);
} else {
this->var_remap_[binding->var] = orig_var;
}
@@ -231,7 +231,7 @@ class CheckpointGenerator : private ExprMutator {
if (it != checkpoint_map_.end()) {
return std::make_pair((*it).second, new_value);
}
- auto new_var = builder_->Emit(new_value, var->name_hint + "_cp");
+ auto new_var = builder_->Emit(new_value, var->name + "_cp");
checkpoint_map_.Set(var, new_var);
return std::make_pair(new_var, new_value);
}
@@ -250,7 +250,7 @@ class CheckpointGenerator : private ExprMutator {
if (it != checkpoint_map_.end()) {
return (*it).second;
}
- Var new_var = builder_->Emit(VisitExpr(binding_map_[var]), var->name_hint
+ "_cp");
+ Var new_var = builder_->Emit(VisitExpr(binding_map_[var]), var->name +
"_cp");
checkpoint_map_.Set(var, new_var);
return new_var;
}
@@ -520,9 +520,9 @@ class BackwardBindingGenerator : private ExprVisitor {
Var EmitAdjoint(const Var& source_var, const Expr& adjoint, bool is_output) {
Var adjoint_var;
if (is_output) {
- adjoint_var = builder_->EmitOutput(adjoint, source_var->name_hint +
"_adjoint_out");
+ adjoint_var = builder_->EmitOutput(adjoint, source_var->name +
"_adjoint_out");
} else {
- adjoint_var = builder_->Emit(adjoint, source_var->name_hint +
"_adjoint");
+ adjoint_var = builder_->Emit(adjoint, source_var->name + "_adjoint");
adjoint_var_map_.Set(source_var, adjoint_var);
}
return adjoint_var;
@@ -768,16 +768,15 @@ class GradientMutator : private ExprMutator {
for (const auto& var : require_grads) {
auto it = var_map.find(var);
TVM_FFI_ICHECK(it != var_map.end())
- << "There is no Var named " << var->name_hint << " in the function "
<< func_name;
- TVM_FFI_ICHECK_EQ(var_set.count(var), 0)
- << "Var " << var->name_hint << " appears more than once";
+ << "There is no Var named " << var->name << " in the function " <<
func_name;
+ TVM_FFI_ICHECK_EQ(var_set.count(var), 0) << "Var " << var->name << "
appears more than once";
var_set.emplace(var);
mapped_vars.push_back((*it).second);
TVM_FFI_ICHECK(IsNestedTensorConditioned(GetType(var),
IsFloatTensorType))
<< "Only Tensors of floating point dtype or Tuples of float "
"Tensors can require gradients, but the Type of Var "
- << var->name_hint << " is " << GetType(var);
+ << var->name << " is " << GetType(var);
}
return mapped_vars;
}
diff --git a/src/relax/transform/inline_functions.cc
b/src/relax/transform/inline_functions.cc
index a15d896457..d790737adb 100644
--- a/src/relax/transform/inline_functions.cc
+++ b/src/relax/transform/inline_functions.cc
@@ -129,7 +129,7 @@ class FunctionInliner : public ExprMutator {
//
// This implementation uses Option 4.
- Var param_var(func->params[i]->name_hint, args[i]->ty.as<Type>());
+ Var param_var(func->params[i]->name, args[i]->ty.as<Type>());
param_bindings.push_back(VarBinding(param_var, args[i]));
param_map.Set(func->params[i], param_var);
}
diff --git a/src/relax/transform/lambda_lift.cc
b/src/relax/transform/lambda_lift.cc
index 6a418ec815..350bb66c90 100644
--- a/src/relax/transform/lambda_lift.cc
+++ b/src/relax/transform/lambda_lift.cc
@@ -96,7 +96,7 @@ class LambdaNameCollector : ExprVisitor {
lifted_with_global_symbol_.insert({func, public_name});
}
- name_stack_.push_back(binding->var->name_hint);
+ name_stack_.push_back(binding->var->name);
lambda_location_.insert({func, name_stack_});
ExprVisitor::VisitBinding_(binding, func);
name_stack_.pop_back();
@@ -285,7 +285,7 @@ class LambdaLifter : public ExprMutator {
ffi::Array<Var> typed_captured_vars;
ffi::Map<Var, Expr> rebinding_map;
for (auto free_var : captured_vars) {
- Var var = Var(free_var->name_hint, GetType(free_var), free_var->span);
+ Var var = Var(free_var->name, GetType(free_var), free_var->span);
typed_captured_vars.push_back(var);
rebinding_map.Set(free_var, var);
}
diff --git a/src/relax/transform/lazy_transform_params.cc
b/src/relax/transform/lazy_transform_params.cc
index 027d796245..e2ce85896b 100644
--- a/src/relax/transform/lazy_transform_params.cc
+++ b/src/relax/transform/lazy_transform_params.cc
@@ -99,10 +99,10 @@ class LazyInputMutator : public ExprMutator {
auto untyped = builder_->Emit(Call(Type::Missing(), plan_->fget_param,
{
PrimExpr(IntImm::Int64(it->second)),
- StringImm(var->name_hint),
+ StringImm(var->name),
}),
- var->name_hint + "_untyped");
- return builder_->EmitMatchCast(untyped, GetType(var), var->name_hint);
+ var->name + "_untyped");
+ return builder_->EmitMatchCast(untyped, GetType(var), var->name);
}
}
diff --git a/src/relax/transform/lift_transform_params.cc
b/src/relax/transform/lift_transform_params.cc
index 0a4214b878..84e77c6394 100644
--- a/src/relax/transform/lift_transform_params.cc
+++ b/src/relax/transform/lift_transform_params.cc
@@ -104,7 +104,7 @@ struct BaseCollectInfo {
}
for (const auto& var : outputs) {
- Var out_var(var->name_hint + "_output", GetType(var));
+ Var out_var(var->name + "_output", GetType(var));
output_var_binding.push_back(VarBinding(out_var, var));
output_exprs.push_back(out_var);
}
@@ -274,7 +274,7 @@ struct LocalCollectInfo : public BaseCollectInfo {
return global_outputs;
}();
for (const auto& var : compile_time_outputs) {
- Var param_var(var->name_hint, GetType(var));
+ Var param_var(var->name, GetType(var));
bindings.push_back(VarBinding(var, param_var));
params.push_back(param_var);
}
diff --git a/src/relax/transform/realize_vdevice.cc
b/src/relax/transform/realize_vdevice.cc
index d2ff9aa5d1..834fe97351 100644
--- a/src/relax/transform/realize_vdevice.cc
+++ b/src/relax/transform/realize_vdevice.cc
@@ -368,9 +368,9 @@ class VDeviceTypeUpdater : ExprMutator {
}();
if (var->IsInstance<DataflowVarNode>()) {
- var = DataflowVar(var->name_hint, new_ty, var->span);
+ var = DataflowVar(var->name, new_ty, var->span);
} else {
- var = Var(var->name_hint, new_ty, var->span);
+ var = Var(var->name, new_ty, var->span);
}
}
}
diff --git a/src/relax/transform/remove_unused_parameters.cc
b/src/relax/transform/remove_unused_parameters.cc
index 61ead90cf7..0ec73ee2d6 100644
--- a/src/relax/transform/remove_unused_parameters.cc
+++ b/src/relax/transform/remove_unused_parameters.cc
@@ -103,7 +103,7 @@ std::optional<CalleeAnalysis> AnalyzeCallee(Function func) {
// Promote the free symbolic var via a 1-D shape param so the param
actually
// *defines* the var. A PrimType param only carries a dtype and defines no
// TIR var, which leaves the var undefined under the strict tirx verifier.
- Var relax_var("param_" + tir_var->name_hint,
ShapeType({tir_var.as_or_throw<PrimExpr>()}));
+ Var relax_var("param_" + tir_var->name,
ShapeType({tir_var.as_or_throw<PrimExpr>()}));
params.push_back(relax_var);
}
diff --git a/src/relax/transform/rewrite_cuda_graph.cc
b/src/relax/transform/rewrite_cuda_graph.cc
index c993a83748..5396d3ce2e 100644
--- a/src/relax/transform/rewrite_cuda_graph.cc
+++ b/src/relax/transform/rewrite_cuda_graph.cc
@@ -138,7 +138,7 @@ class FuncBuilder : public ExprMutator {
}
// Set up the parameters
for (const auto* input : inputs_) {
- auto new_var = Var(input->name_hint,
VisitExprDepTypeField(input->ty.as_or_throw<Type>()));
+ auto new_var = Var(input->name,
VisitExprDepTypeField(input->ty.as_or_throw<Type>()));
var_remap_[ffi::GetRef<Var>(input)] = new_var;
params.push_back(new_var);
}
@@ -252,7 +252,7 @@ class CUDAGraphRewritePlanner : public ExprVisitor {
if (i < num_inputs) {
for (const auto& symbolic_var : symbolic_vars) {
auto prim_var = symbolic_var.as_or_throw<tirx::PrimVar>();
- if
(capture_symbolic_var_name_hints.count(symbolic_var->name_hint)) {
+ if (capture_symbolic_var_name_hints.count(symbolic_var->name)) {
capture_symbolic_vars_.insert(prim_var);
}
}
@@ -869,7 +869,7 @@ class CUDAGraphRewriter : public ExprMutator {
}
Var EmitRedef(const VarNode* var, const Expr& redef) {
- auto new_var = builder_->Emit(redef, var->name_hint);
+ auto new_var = builder_->Emit(redef, var->name);
var_remap_[ffi::GetRef<Var>(var)] = new_var;
return new_var;
}
diff --git a/src/relax/transform/specialize_primfunc_based_on_callsite.cc
b/src/relax/transform/specialize_primfunc_based_on_callsite.cc
index e9563d43cd..9425694102 100644
--- a/src/relax/transform/specialize_primfunc_based_on_callsite.cc
+++ b/src/relax/transform/specialize_primfunc_based_on_callsite.cc
@@ -94,7 +94,7 @@ class SpecializeTIRCallArgs : ExprMutator {
}
ffi::String name;
if (args[i]->IsInstance<tvm::VarNode>()) {
- name = args[i].as_or_throw<Var>()->name_hint;
+ name = args[i].as_or_throw<Var>()->name;
} else {
name = std::string({static_cast<char>('A' + i)});
}
diff --git a/src/relax/transform/static_plan_block_memory.cc
b/src/relax/transform/static_plan_block_memory.cc
index c5db48b532..7d20152ab2 100644
--- a/src/relax/transform/static_plan_block_memory.cc
+++ b/src/relax/transform/static_plan_block_memory.cc
@@ -442,8 +442,8 @@ void SetTIRVarRangeConstraints(Function func,
arith::AnalyzerObj* ana,
}
ffi::Array<tirx::Var> var_in_signature = TIRVarsInType(GetType(func));
for (const tirx::Var& tir_var : var_in_signature) {
- auto it_upper = var_upper_bound_attr.find(tir_var->name_hint);
- auto it_lower = var_lower_bound_attr.find(tir_var->name_hint);
+ auto it_upper = var_upper_bound_attr.find(tir_var->name);
+ auto it_lower = var_lower_bound_attr.find(tir_var->name);
// Only bind the variable to a range if an upper bound is explicitly
provided.
// Without an upper bound, memory planning cannot determine the required
storage size,
@@ -457,7 +457,7 @@ void SetTIRVarRangeConstraints(Function func,
arith::AnalyzerObj* ana,
dom_map->Set(tir_var, arith::IntSet::FromRange(range));
} else if (it_lower != var_lower_bound_attr.end() &&
it_lower->second->value >= 0) {
ana->MarkGlobalNonNegValue(tir_var.as_or_throw<PrimExpr>());
- } else if (non_negative_var_attr.count(tir_var->name_hint)) {
+ } else if (non_negative_var_attr.count(tir_var->name)) {
ana->MarkGlobalNonNegValue(tir_var.as_or_throw<PrimExpr>());
}
}
diff --git a/src/relax/transform/to_mixed_precision.cc
b/src/relax/transform/to_mixed_precision.cc
index 0764fd268f..3fe3da71b7 100644
--- a/src/relax/transform/to_mixed_precision.cc
+++ b/src/relax/transform/to_mixed_precision.cc
@@ -283,7 +283,7 @@ class ToMixedPrecisionRewriter : public ExprMutator {
if (it != var_remap_.end()) {
return it->second;
} else {
- if (fp16_input_names_.count(var->name_hint)) {
+ if (fp16_input_names_.count(var->name)) {
auto ty = GetType(var);
if (auto tensor_ty = ty.as<TensorTypeNode>()) {
VDevice vdev = VDevice();
@@ -291,7 +291,7 @@ class ToMixedPrecisionRewriter : public ExprMutator {
vdev = tensor_ty->vdevice.value();
}
TensorType fp16_ty(tensor_ty->shape.value(), PrimType::Float(16),
vdev, tensor_ty->span);
- Var fp16_var(var->name_hint, fp16_ty, var->span);
+ Var fp16_var(var->name, fp16_ty, var->span);
var_remap_[var] = fp16_var;
return fp16_var;
}
diff --git a/src/relax/transform/to_non_dataflow.cc
b/src/relax/transform/to_non_dataflow.cc
index ae279e4cd5..695e700251 100644
--- a/src/relax/transform/to_non_dataflow.cc
+++ b/src/relax/transform/to_non_dataflow.cc
@@ -34,7 +34,7 @@ class ToNonDFMutator : public ExprMutator {
public:
Var VisitVarDef(const Var& var) final {
if (var.as<DataflowVarNode>()) {
- Var new_var = Var(var->name_hint, GetType(var), var->span);
+ Var new_var = Var(var->name, GetType(var), var->span);
this->var_remap_[var] = new_var;
return new_var;
}
diff --git a/src/relax/transform/utils.h b/src/relax/transform/utils.h
index 54b6d9f5a0..2f2e0895e7 100644
--- a/src/relax/transform/utils.h
+++ b/src/relax/transform/utils.h
@@ -245,9 +245,9 @@ class SymbolicVarRenewMutator : public ExprMutator {
static Var CopyVar(const VarNode* op, Type ty) {
ffi::Optional<Type> ty_annotation = ty.IsMissing() ? std::nullopt :
ffi::Optional<Type>(ty);
if (op->IsInstance<DataflowVarNode>()) {
- return DataflowVar(op->name_hint, std::move(ty_annotation), op->span);
+ return DataflowVar(op->name, std::move(ty_annotation), op->span);
}
- return Var(op->name_hint, std::move(ty_annotation), op->span);
+ return Var(op->name, std::move(ty_annotation), op->span);
}
Type RenewType(const VarNode* op) {
diff --git a/src/s_tir/analysis/verify_gpu_code.cc
b/src/s_tir/analysis/verify_gpu_code.cc
index c198a1cd06..9f20b80a0e 100644
--- a/src/s_tir/analysis/verify_gpu_code.cc
+++ b/src/s_tir/analysis/verify_gpu_code.cc
@@ -107,7 +107,7 @@ class GPUCodeVerifier : public StmtExprVisitor {
const auto* extent = op->value.as<IntImmNode>();
TVM_FFI_ICHECK(extent);
- std::string name = var.get()->name_hint;
+ std::string name = var.get()->name;
// record the number of threads in a block
if (name == "threadIdx.x" || name == "threadIdx.y" || name ==
"threadIdx.z" ||
name == "vthread") {
@@ -183,7 +183,7 @@ class GPUCodeVerifier : public StmtExprVisitor {
}
void VisitStmt_(const ForNode* op) {
- if (op->loop_var->name_hint == "vthread.s") {
+ if (op->loop_var->name == "vthread.s") {
const auto* extent = op->extent.as<IntImmNode>();
TVM_FFI_ICHECK(extent);
diff --git a/src/s_tir/data_layout.cc b/src/s_tir/data_layout.cc
index e449208904..4393bfabb3 100644
--- a/src/s_tir/data_layout.cc
+++ b/src/s_tir/data_layout.cc
@@ -71,7 +71,7 @@ const SLayoutAxis& SLayoutAxis::Get(const char name) {
}
const SLayoutAxis& SLayoutAxis::Get(const IterVar& itvar) {
- const std::string axis = itvar->var.get()->name_hint;
+ const std::string axis = itvar->var.get()->name;
TVM_FFI_ICHECK_EQ(axis.size(), 1) << "Invalid layout axis " << axis;
return SLayoutAxis::Get(axis[0]);
}
@@ -99,12 +99,12 @@ SLayout::SLayout(const ffi::Array<IterVar>& axes) {
TVM_FFI_ICHECK(!is_grouped)
<< "Only Subordinate Axes with extent is allowed within a packed
dim";
}
- TVM_FFI_ICHECK_EQ(axis->var.get()->name_hint.size(), 1)
- << "Invalid layout axis " << axis->var.get()->name_hint;
- char c = axis->var.get()->name_hint.operator std::string()[0];
+ TVM_FFI_ICHECK_EQ(axis->var.get()->name.size(), 1)
+ << "Invalid layout axis " << axis->var.get()->name;
+ char c = axis->var.get()->name.operator std::string()[0];
TVM_FFI_ICHECK((c >= 'a' && c <= 'z') || (c >= 'A' && c <= 'Z'))
<< "Invalid layout axis " << c;
- repr << axis->var.get()->name_hint;
+ repr << axis->var.get()->name;
}
if (is_grouped) repr << "]";
}
@@ -167,9 +167,9 @@ SLayout::SLayout(const std::string& name, PrimType
index_ty) { // NOLINT(*)
int64_t extent = 1;
for (auto& axis : unpacked_axes) {
TVM_FFI_ICHECK(axis->dom->extent.as<IntImmNode>())
- << "Invalid SLayout " << name << ": can't have variable sized
node("
- << axis->var->name_hint << ") within a packed axis";
- auto axis_name = axis->var->name_hint.operator std::string();
+ << "Invalid SLayout " << name << ": can't have variable sized
node(" << axis->var->name
+ << ") within a packed axis";
+ auto axis_name = axis->var->name.operator std::string();
auto factor = axis->dom->extent.as<IntImm>().value();
ss << axis_name;
extent = extent * factor->value;
@@ -192,7 +192,7 @@ SLayout::SLayout(const std::string& name, PrimType
index_ty) { // NOLINT(*)
std::vector<int> axis_cnt(256, 0);
for (const IterVar& pv : node->axes) {
for (const IterVar& v : UnpackIterVar(pv)) {
- auto axis_str = v->var.get()->name_hint.operator std::string();
+ auto axis_str = v->var.get()->name.operator std::string();
TVM_FFI_ICHECK_EQ(axis_str.size(), 1);
char axis = axis_str[0];
TVM_FFI_ICHECK((axis >= 'a' && axis <= 'z') || (axis >= 'A' && axis <=
'Z'));
@@ -201,7 +201,7 @@ SLayout::SLayout(const std::string& name, PrimType
index_ty) { // NOLINT(*)
}
for (const IterVar& pv : node->axes) {
for (const IterVar& v : UnpackIterVar(pv)) {
- char axis = v->var.get()->name_hint.operator std::string()[0];
+ char axis = v->var.get()->name.operator std::string()[0];
if (axis >= 'a' && axis <= 'z') {
TVM_FFI_ICHECK(axis_cnt[axis - 'a' + 'A'])
<< "Invalid layout " << name << ": missing axis " <<
std::toupper(axis);
@@ -231,7 +231,7 @@ ffi::Array<IterVar> SLayout::UnpackIterVar(IterVar
packed_iter) {
ffi::Array<IterVar> result;
int64_t factor = 0, final_factor = 1;
- std::string name(packed_iter->var->name_hint.c_str());
+ std::string name(packed_iter->var->name.c_str());
PrimType index_ty = packed_iter->var.ty();
for (auto ch : name) {
@@ -261,7 +261,7 @@ IterVar SLayout::PackIterVar(ffi::Array<IterVar> iter_vars)
{
for (auto itvar : iter_vars) {
TVM_FFI_ICHECK(itvar->dom->extent.as<IntImm>())
<< "Packed Axis can contain only Subordinate Axes";
- name << itvar->dom->extent.as<IntImm>().value() << itvar->var->name_hint;
+ name << itvar->dom->extent.as<IntImm>().value() << itvar->var->name;
extent = extent * itvar->dom->extent.as<IntImm>().value()->value;
}
@@ -594,7 +594,7 @@ TVM_FFI_STATIC_INIT_BLOCK() {
.def("s_tir.SLayoutGetItem",
[](SLayout layout, int idx) -> std::string {
const auto& axis = layout.PackedAxisAt(idx);
- return axis->var->name_hint;
+ return axis->var->name;
})
.def("s_tir.SBijectiveLayout",
[](SLayout src_layout, SLayout dst_layout) -> SBijectiveLayout {
diff --git a/src/s_tir/schedule/primitive/annotate_buffer_access.cc
b/src/s_tir/schedule/primitive/annotate_buffer_access.cc
index 10887e2c4d..0e2448f87c 100644
--- a/src/s_tir/schedule/primitive/annotate_buffer_access.cc
+++ b/src/s_tir/schedule/primitive/annotate_buffer_access.cc
@@ -134,7 +134,7 @@ struct AnnotateBufferAccessTraits : public
UnpackedInstTraits<AnnotateBufferAcce
std::ostringstream oss;
auto print_expr = [&oss](const PrimExpr& expr) {
if (auto var = expr.as<PrimVar>()) {
- oss << var.value()->name_hint;
+ oss << var.value()->name;
} else {
oss << expr;
}
@@ -142,7 +142,7 @@ struct AnnotateBufferAccessTraits : public
UnpackedInstTraits<AnnotateBufferAcce
oss << "lambda ";
for (size_t i = 0; i < index_map->initial_indices.size(); ++i) {
if (i != 0) oss << ", ";
- oss << index_map->initial_indices[i]->name_hint;
+ oss << index_map->initial_indices[i]->name;
}
oss << ": [";
for (size_t i = 0; i < index_map->final_indices.size(); i += 2) {
diff --git a/src/s_tir/schedule/primitive/cache_index.cc
b/src/s_tir/schedule/primitive/cache_index.cc
index 9871f18fd3..e0c075544f 100644
--- a/src/s_tir/schedule/primitive/cache_index.cc
+++ b/src/s_tir/schedule/primitive/cache_index.cc
@@ -270,7 +270,7 @@ ffi::Array<SBlock> MakeIndexCacheStage(IndexInfo* info,
const ffi::String& stora
arith::EvalSet(info->var_binding.at(it),
arith::AsIntSet(info->range_map)).max() + 1);
}
info->cache_buffer.push_back(Buffer(index_buffer_var, data_ty,
buffer_shape, {1}, {0},
- index_buffer_var->name_hint, 0, 0,
kDefault));
+ index_buffer_var->name, 0, 0,
kDefault));
// Create loop vars and block vars' binding_value
std::vector<PrimVar> loop_vars;
diff --git a/src/s_tir/schedule/primitive/cache_read_write.cc
b/src/s_tir/schedule/primitive/cache_read_write.cc
index 47aa8354a7..c1d2025e9b 100644
--- a/src/s_tir/schedule/primitive/cache_read_write.cc
+++ b/src/s_tir/schedule/primitive/cache_read_write.cc
@@ -165,7 +165,7 @@ SBlock MakeReindexCacheStage(const BufferRegion&
cache_region, ReindexCacheStage
ffi::Map<Var, Var> var_map;
for (size_t i = 0; i < info->loop_vars.size(); ++i) {
Var original_var = info->loop_vars[i];
- PrimVar loop_var(original_var->name_hint,
original_var->ty.as_or_throw<PrimType>());
+ PrimVar loop_var(original_var->name,
original_var->ty.as_or_throw<PrimType>());
var_map.Set(original_var, loop_var);
loop_vars.push_back(loop_var);
}
@@ -174,7 +174,7 @@ SBlock MakeReindexCacheStage(const BufferRegion&
cache_region, ReindexCacheStage
PrimExpr original_iter_value = info->block_iter_values[i];
IterVar block_var = IterVar(
/*dom=*/original_block_var->dom,
- /*var=*/PrimVar(original_block_var->var->name_hint,
original_block_var->var.ty()),
+ /*var=*/PrimVar(original_block_var->var->name,
original_block_var->var.ty()),
/*IterVarType=*/kDataPar);
var_map.Set(original_block_var->var, block_var->var);
block_vars.push_back(block_var);
@@ -2039,7 +2039,7 @@ void CollectReindexCacheStageInfoAndCreateBuffer(
ffi::ObjectPtr<VarNode> new_var =
ffi::make_object<VarNode>(*old_buffer->data.get());
const auto* ptr_type = TVM_TYPE_AS(old_buffer->data->ty, PointerTypeNode);
new_var->ty = PointerType(ptr_type->element_type, storage_scope);
- new_buffer->data = Var(new_var->name_hint + "_" + storage_scope,
new_var->ty);
+ new_buffer->data = Var(new_var->name + "_" + storage_scope, new_var->ty);
new_buffer->name = old_buffer->name + "_" + storage_scope;
new_buffer->shape = new_shape;
diff --git a/src/s_tir/schedule/primitive/layout_transformation.cc
b/src/s_tir/schedule/primitive/layout_transformation.cc
index de0ed83e28..9ebdd599a2 100644
--- a/src/s_tir/schedule/primitive/layout_transformation.cc
+++ b/src/s_tir/schedule/primitive/layout_transformation.cc
@@ -294,7 +294,7 @@ class TransformLayoutPlanner : private StmtExprVisitor {
new_indices = inverse->initial_indices.Map([](PrimVar var) {
std::stringstream ss;
- ss << "v_" << var->name_hint;
+ ss << "v_" << var->name;
return Var(ss.str(), var.ty());
});
@@ -478,7 +478,7 @@ class TransformLayoutPlanner : private StmtExprVisitor {
for (size_t i = 0; i < inverse->initial_indices.size(); i++) {
const auto& loop_var = inverse->initial_indices[i];
const auto& dim = new_buffer->shape[i];
- Var block_var("v_" + loop_var->name_hint, loop_var.ty());
+ Var block_var("v_" + loop_var->name, loop_var.ty());
IterVar iter_var(Range(0, dim), block_var.as_or_throw<PrimVar>(),
kDataPar);
loop_indices_to_block_indices.Set(loop_var, block_var);
indices.push_back(iter_var->var);
@@ -573,7 +573,7 @@ class TransformLayoutPlanner : private StmtExprVisitor {
for (size_t i = 0; i < inverse->initial_indices.size(); i++) {
const auto& loop_var = inverse->initial_indices[i];
const auto& dim = new_buffer->shape[i];
- Var block_var("v_" + loop_var->name_hint, loop_var.ty());
+ Var block_var("v_" + loop_var->name, loop_var.ty());
IterVar iter_var(Range(0, dim), block_var.as_or_throw<PrimVar>(),
kDataPar);
indices.push_back(iter_var->var);
iter_vars.push_back(iter_var);
@@ -1135,7 +1135,7 @@ IndexMap LegalizeIndexMapDType(const IndexMap& index_map,
const ffi::Array<PrimE
DLDataType initial_dtype = initial_indices_orig[i].ty()->dtype;
if (arg_dtype != initial_dtype) {
- auto new_idx = Var(initial_indices_orig[i]->name_hint, args[i].ty());
+ auto new_idx = Var(initial_indices_orig[i]->name, args[i].ty());
initial_indices.push_back(new_idx);
var_map.Set(initial_indices_orig[i], new_idx.as_or_throw<PrimExpr>());
} else {
diff --git a/src/s_tir/schedule/primitive/loop_transformation.cc
b/src/s_tir/schedule/primitive/loop_transformation.cc
index b3573843b4..845e88fe64 100644
--- a/src/s_tir/schedule/primitive/loop_transformation.cc
+++ b/src/s_tir/schedule/primitive/loop_transformation.cc
@@ -485,7 +485,7 @@ class BufferIndicesMapExtractor : public StmtExprVisitor {
check_ = true;
break;
}
- indices.push_back(var.value()->name_hint);
+ indices.push_back(var.value()->name);
}
if (buffer_indices_map.find(store->buffer->name) ==
buffer_indices_map.end() && !check_)
buffer_indices_map.Set(store->buffer->name, indices);
@@ -501,7 +501,7 @@ class BufferIndicesMapExtractor : public StmtExprVisitor {
check_ = true;
break;
}
- indices.push_back(var.value()->name_hint);
+ indices.push_back(var.value()->name);
}
if (buffer_indices_map.find(load->buffer->name) ==
buffer_indices_map.end() && !check_)
buffer_indices_map.Set(load->buffer->name, indices);
@@ -547,7 +547,7 @@ class BlockMutator : public StmtExprMutator {
// If iter_vars.size() is 0, then the block most probably be an Opaque
block
if (new_block->iter_vars.size() == 0 || inner_iter_var_index == -1) {
new_block.CopyOnWrite()->name_hint =
- new_block.CopyOnWrite()->name_hint + "_" + new_loop_var_->name_hint;
+ new_block.CopyOnWrite()->name_hint + "_" + new_loop_var_->name;
return new_block;
}
@@ -558,7 +558,7 @@ class BlockMutator : public StmtExprMutator {
MutateArray(new_block->iter_vars, [this, &iter_var_](const IterVar&
iter) {
auto dtype = iter->var.ty();
// Create new Var instance for each IterVar
- Var new_var = Var(iter->var->name_hint, iter->var.ty());
+ Var new_var = Var(iter->var->name, iter->var.ty());
IterVar new_iter = iter;
new_iter.CopyOnWrite()->var = new_var.as_or_throw<PrimVar>();
// Change the domain of IterVar corresponding to partitioned loop_var
@@ -572,14 +572,14 @@ class BlockMutator : public StmtExprMutator {
if (!new_block->iter_vars.same_as(new_iter_vars)) {
new_block.CopyOnWrite()->iter_vars = std::move(new_iter_vars);
new_block.CopyOnWrite()->name_hint =
- new_block.CopyOnWrite()->name_hint + "_" + new_loop_var_->name_hint;
+ new_block.CopyOnWrite()->name_hint + "_" + new_loop_var_->name;
}
// Get the (iter_var, new Range) map
ffi::Map<ffi::String, Range> index_range_map;
for (size_t i = 0; i < new_block->iter_vars.size(); i++) {
IterVar iter = new_block->iter_vars[i];
- index_range_map.Set(iter->var->name_hint, iter->dom);
+ index_range_map.Set(iter->var->name, iter->dom);
}
// Get the (Buffer, indices) map
@@ -623,7 +623,7 @@ class BlockMutator : public StmtExprMutator {
Stmt VisitStmt_(const ForNode* op) final {
For res = StmtMutator::VisitStmt_(op).as_or_throw<For>();
- Var new_var = Var(op->loop_var->name_hint, op->loop_var.ty());
+ Var new_var = Var(op->loop_var->name, op->loop_var.ty());
if (!op->loop_var.same_as(new_var)) {
// If the partioned loop contains nested for loop, then create new
iteration variable instance
@@ -664,7 +664,7 @@ ffi::Array<StmtSRef> LoopPartition(ScheduleState self,
const StmtSRef& loop_sref
dtype = PrimType::Int(bits);
}
- ffi::String block_name = get_sblock_name(loop->body) + "_" +
loop->loop_var->name_hint;
+ ffi::String block_name = get_sblock_name(loop->body) + "_" +
loop->loop_var->name;
int n = factors.size();
PrimExpr min_value = loop->min;
PrimExpr extent_value;
@@ -913,7 +913,7 @@ StmtSRef Fuse(ScheduleState self, const
ffi::Array<StmtSRef>& loop_srefs,
return false;
};
if (UsesVar(loop->extent, f_contain)) {
- throw DependentLoopError(self->mod, ffi::GetRef<For>(loop),
used_var->name_hint,
+ throw DependentLoopError(self->mod, ffi::GetRef<For>(loop),
used_var->name,
DependentLoopError::PrimitiveKind::kFuse);
}
outer_loop_vars.insert(loop->loop_var.get());
@@ -924,7 +924,7 @@ StmtSRef Fuse(ScheduleState self, const
ffi::Array<StmtSRef>& loop_srefs,
int n = loops.size();
int bits = loops[0]->loop_var.ty().bits();
for (int i = 1; i < n; i++) {
- suffix += "_" + loops[i]->loop_var->name_hint;
+ suffix += "_" + loops[i]->loop_var->name;
bits = std::max(bits, loops[i]->loop_var.ty().bits());
}
suffix += "_fused";
@@ -1111,7 +1111,7 @@ For ConstructNewLoopChain(const ScheduleState& self,
std::vector<const StmtSRefN
return false;
};
if (UsesVar(copy->min, f_contain) || UsesVar(copy->extent, f_contain)) {
- throw DependentLoopError(self->mod, ffi::GetRef<For>(copy),
used_var->name_hint,
+ throw DependentLoopError(self->mod, ffi::GetRef<For>(copy),
used_var->name,
DependentLoopError::PrimitiveKind::kReorder);
}
inner_vars.insert(copy->loop_var.get());
diff --git a/src/s_tir/schedule/primitive/reduction.cc
b/src/s_tir/schedule/primitive/reduction.cc
index 9dbc58233e..cf5e74a51e 100644
--- a/src/s_tir/schedule/primitive/reduction.cc
+++ b/src/s_tir/schedule/primitive/reduction.cc
@@ -889,7 +889,7 @@ class RFactorBlockCreator : public BaseBlockCreator {
void CreateAdditionalIter() final {
// Create a new data parallel block iter for the rfactor loop.
additional_iter_ =
- IterVarFromLoop(rf_loop_, "v" + rf_loop_->loop_var->name_hint,
IterVarType::kDataPar);
+ IterVarFromLoop(rf_loop_, "v" + rf_loop_->loop_var->name,
IterVarType::kDataPar);
loop_var2block_binding_[rf_loop_->loop_var.get()] = additional_iter_->var;
iter_vars_.push_back(additional_iter_);
iter_values_.push_back(rf_loop_->loop_var);
@@ -923,7 +923,7 @@ class RFactorBlockCreator : public BaseBlockCreator {
// We haven't created the new block iter for `var`. So here we create
it, append it
// and its binding to `rf_block_iter_vars` and `rf_block_iter_values`
respectively.
IterVar new_iter_var =
- IterVarFromLoop(loop, "v" + loop->loop_var->name_hint,
IterVarType::kCommReduce);
+ IterVarFromLoop(loop, "v" + loop->loop_var->name,
IterVarType::kCommReduce);
loop_var2block_binding_[var.get()] = new_iter_var->var;
iter_vars_.push_back(new_iter_var);
iter_values_.push_back(var.as_or_throw<PrimExpr>());
@@ -1017,7 +1017,7 @@ class WriteBackBlockCreator : public BaseBlockCreator {
void CreateAdditionalIter() final {
// Create a new reduction block iter for the rfactor loop.
IterVar wb_new_block_iter =
- IterVarFromLoop(rf_loop_, "v" + rf_loop_->loop_var->name_hint,
kCommReduce);
+ IterVarFromLoop(rf_loop_, "v" + rf_loop_->loop_var->name, kCommReduce);
iter_vars_.push_back(wb_new_block_iter);
iter_values_.push_back(rf_loop_->loop_var);
var_map_.Set(rf_additional_iter_->var, wb_new_block_iter->var);
diff --git a/src/s_tir/schedule/trace.cc b/src/s_tir/schedule/trace.cc
index 6410416baa..516a44d287 100644
--- a/src/s_tir/schedule/trace.cc
+++ b/src/s_tir/schedule/trace.cc
@@ -205,7 +205,7 @@ ffi::Array<Any> TranslateInputRVs(
if (obj.as<IndexMapNode>()) {
IndexMap index_map = obj.as_or_throw<IndexMap>();
index_map = Substitute(index_map, [&named_rvs](const Var& var) ->
ffi::Optional<PrimExpr> {
- auto it = named_rvs.find(var->name_hint);
+ auto it = named_rvs.find(var->name);
if (it != named_rvs.end()) {
return it->second.as_or_throw<Var>().as_or_throw<PrimExpr>();
}
diff --git a/src/s_tir/schedule/transform.cc b/src/s_tir/schedule/transform.cc
index ea67a44ce8..4c7ed8a074 100644
--- a/src/s_tir/schedule/transform.cc
+++ b/src/s_tir/schedule/transform.cc
@@ -43,7 +43,7 @@ Buffer WithScope(const Buffer& buffer, const ffi::String&
scope) {
ffi::ObjectPtr<BufferNode> new_buffer =
ffi::make_object<BufferNode>(*buffer.get());
const auto* ptr_type = TVM_TYPE_AS(buffer->data->ty, PointerTypeNode);
Type new_type = PointerType(ptr_type->element_type, scope);
- new_buffer->data = tirx::Var(buffer->data->name_hint + "_" + scope,
new_type);
+ new_buffer->data = tirx::Var(buffer->data->name + "_" + scope, new_type);
new_buffer->name = buffer->name + "_" + scope;
return Buffer(new_buffer);
}
@@ -52,8 +52,7 @@ Buffer WithDType(const Buffer& buffer, PrimType dtype) {
ffi::ObjectPtr<BufferNode> new_buffer =
ffi::make_object<BufferNode>(*buffer.get());
new_buffer->dtype = dtype;
const auto* ptr_type = TVM_TYPE_AS(buffer->data->ty, PointerTypeNode);
- new_buffer->data =
- tirx::Var(buffer->data->name_hint, PointerType(dtype,
ptr_type->storage_scope));
+ new_buffer->data = tirx::Var(buffer->data->name, PointerType(dtype,
ptr_type->storage_scope));
new_buffer->name = buffer->name;
return Buffer(new_buffer);
}
diff --git a/src/s_tir/transform/compact_buffer_region.cc
b/src/s_tir/transform/compact_buffer_region.cc
index 2db7d08278..b963af8b40 100644
--- a/src/s_tir/transform/compact_buffer_region.cc
+++ b/src/s_tir/transform/compact_buffer_region.cc
@@ -362,7 +362,7 @@ class BufferAccessRegionCollector : public StmtExprVisitor {
}
auto dom_it = dom_map_.find(v);
TVM_FFI_ICHECK(dom_it != dom_map_.end())
- << "Could not find domain for loop variable " << v->name_hint;
+ << "Could not find domain for loop variable " << v->name;
non_relaxed[i] = dom_it->second;
dom_map_.erase(dom_it);
}
diff --git a/src/s_tir/transform/convert_blocks_to_opaque.cc
b/src/s_tir/transform/convert_blocks_to_opaque.cc
index 3bc9e9e514..15c7b2b06b 100644
--- a/src/s_tir/transform/convert_blocks_to_opaque.cc
+++ b/src/s_tir/transform/convert_blocks_to_opaque.cc
@@ -49,7 +49,7 @@ class OpaqueBlockConverter : public StmtExprMutator {
Expr VisitExpr_(const VarNode* var) final {
TVM_FFI_ICHECK(!forbidden_iter_vars_.count(var))
- << "Variable " << var->name_hint << " occurs in the predicate or
iter_values of a block, "
+ << "Variable " << var->name << " occurs in the predicate or
iter_values of a block, "
<< "but isn't defined until the body of the block";
auto it = var_substitutes_.find(var);
diff --git a/src/s_tir/transform/inject_double_buffer.cc
b/src/s_tir/transform/inject_double_buffer.cc
index 151bc11f55..d01ce9e6b4 100644
--- a/src/s_tir/transform/inject_double_buffer.cc
+++ b/src/s_tir/transform/inject_double_buffer.cc
@@ -168,7 +168,7 @@ class DoubleBufferInjector : public StmtExprMutator {
PrimExpr factor = IntImm(new_ext.ty(), split_loop_);
PrimExpr outer_ext = new_ext / factor;
PrimExpr tail_base = outer_ext * factor;
- Var outer_var(old_loop->loop_var->name_hint + ".outer",
old_loop->loop_var.ty());
+ Var outer_var(old_loop->loop_var->name + ".outer",
old_loop->loop_var.ty());
std::unordered_map<const VarNode*, PrimExpr> vmap;
std::vector<Stmt> loop_seq;
for (int32_t i = 0; i < split_loop_; ++i) {
@@ -280,7 +280,7 @@ class DoubleBufferInjector : public StmtExprMutator {
PrimExpr one = IntImm(e.loop->loop_var.ty(), 1);
PrimExpr two = IntImm(e.loop->loop_var.ty(), 2);
PrimExpr loop_shift = e.loop->loop_var + one;
- e.switch_write_var = Var(e.loop->loop_var->name_hint + ".db",
e.loop->loop_var.ty());
+ e.switch_write_var = Var(e.loop->loop_var->name + ".db",
e.loop->loop_var.ty());
e.switch_read_var = indexmod(e.loop->loop_var, two);
in_double_buffer_scope_ = true;
Stmt body = this->VisitStmt(op->body);
diff --git a/src/s_tir/transform/inject_virtual_thread.cc
b/src/s_tir/transform/inject_virtual_thread.cc
index b1740260f3..e542707028 100644
--- a/src/s_tir/transform/inject_virtual_thread.cc
+++ b/src/s_tir/transform/inject_virtual_thread.cc
@@ -481,7 +481,7 @@ class VTInjector : public arith::IRMutatorWithAnalyzer {
return SeqStmt::Flatten(seq);
} else {
// insert a for loop
- Var idx(var_->name_hint + ".s", var_->ty);
+ Var idx(var_->name + ".s", var_->ty);
stmt = Substitute(stmt, ffi::Map<Var, Expr>{{var_, idx}});
PrimType idx_dtype = idx->ty.as_or_throw<PrimType>();
return For(idx.as_or_throw<PrimVar>(), IntImm(idx_dtype, 0),
diff --git a/src/s_tir/transform/lower_cross_thread_reduction.cc
b/src/s_tir/transform/lower_cross_thread_reduction.cc
index c1fefb5c2e..23011943b4 100644
--- a/src/s_tir/transform/lower_cross_thread_reduction.cc
+++ b/src/s_tir/transform/lower_cross_thread_reduction.cc
@@ -444,7 +444,7 @@ Stmt TransformReductionBlock(const SBlockRealizeNode*
realize,
IterVar new_iter_var{nullptr};
{
ffi::ObjectPtr<IterVarNode> n =
ffi::make_object<IterVarNode>(*iter_var.get());
- Var v(iter_var->var->name_hint, iter_var->var->ty,
iter_var->var->span);
+ Var v(iter_var->var->name, iter_var->var->ty, iter_var->var->span);
n->var = v.as_or_throw<PrimVar>();
new_iter_var = IterVar(n);
}
@@ -670,7 +670,7 @@ class CrossThreadReductionTransformer : public StmtMutator {
TVM_FFI_CHECK(IsBoundToThreadIdx(reduction_loop), ValueError)
<< "Cross-thread reduction requires all the reduction-related
loops that "
"are bound to GPU thread axes to only be bound
`threadIdx.x/y/z`. However, loop "
- << reduction_loop->loop_var->name_hint << " violates the
condition.";
+ << reduction_loop->loop_var->name << " violates the condition.";
}
}
diff --git a/src/s_tir/transform/memhammer_coalesce.cc
b/src/s_tir/transform/memhammer_coalesce.cc
index 57388f4e81..728bcace44 100644
--- a/src/s_tir/transform/memhammer_coalesce.cc
+++ b/src/s_tir/transform/memhammer_coalesce.cc
@@ -37,7 +37,7 @@ Stmt FuseNestLoops(Stmt body) {
std::string suffix;
int n = loops.size();
for (int i = 1; i < n; i++) {
- suffix += "_" + loops[i]->loop_var->name_hint;
+ suffix += "_" + loops[i]->loop_var->name;
}
suffix += "_fused";
PrimVar fused_var = loops[0]->loop_var.CopyWithSuffix(suffix);
diff --git a/src/s_tir/transform/merge_shared_memory_allocations.cc
b/src/s_tir/transform/merge_shared_memory_allocations.cc
index 1bed17389f..952468f632 100644
--- a/src/s_tir/transform/merge_shared_memory_allocations.cc
+++ b/src/s_tir/transform/merge_shared_memory_allocations.cc
@@ -394,7 +394,7 @@ class SharedMemoryRewriter : public StmtExprMutator {
// 7. Wrap with the merged-buffer AllocBuffer.
Buffer merged_buf(scope.merged_buf_var, PrimType::UInt(8),
{scope.merged_alloc_size}, {},
- PrimExpr(), scope.merged_buf_var->name_hint, 0, 0,
BufferType::kDefault);
+ PrimExpr(), scope.merged_buf_var->name, 0, 0,
BufferType::kDefault);
ffi::Map<ffi::String, ffi::Any> annotations;
if (scope.has_volatile_alloc) {
annotations.Set(tirx::attr::kVolatile, true);
diff --git a/src/s_tir/transform/renew_defs.cc
b/src/s_tir/transform/renew_defs.cc
index 392e89b86e..23410ae422 100644
--- a/src/s_tir/transform/renew_defs.cc
+++ b/src/s_tir/transform/renew_defs.cc
@@ -156,7 +156,7 @@ class RenewDefMutator : public StmtExprMutator {
private:
Var ReDefineVar(const Var& var) {
- Var new_var(var->name_hint, var->ty, var->span);
+ Var new_var(var->name, var->ty, var->span);
this->AddDefRemap(var, new_var);
return new_var;
}
diff --git a/src/s_tir/transform/tensorcore_infer_fragment.cc
b/src/s_tir/transform/tensorcore_infer_fragment.cc
index 860cdc2cc0..28b92a6e62 100644
--- a/src/s_tir/transform/tensorcore_infer_fragment.cc
+++ b/src/s_tir/transform/tensorcore_infer_fragment.cc
@@ -163,11 +163,11 @@ class FragmentChecker : public StmtExprVisitor {
// A tool for checking shapes of two fragments
bool CheckShape(const VarNode* buffer1, const VarNode* buffer2) {
TVM_FFI_ICHECK(fragment_getter.fragments.count(buffer1))
- << "Tensorecore fragment " << buffer1->name_hint
+ << "Tensorecore fragment " << buffer1->name
<< " must be filled (with tvm_fill_fragment) or loaded (with
tvm_load_matrix_sync) before "
"use.";
TVM_FFI_ICHECK(fragment_getter.fragments.count(buffer2))
- << "Tensorecore fragment " << buffer2->name_hint
+ << "Tensorecore fragment " << buffer2->name
<< " must be filled (with tvm_fill_fragment) or loaded (with
tvm_load_matrix_sync) before "
"use.";
FragmentInfo info1 = fragment_getter.fragments.at(buffer1);
diff --git a/src/script/ir_builder/ir/ir.cc b/src/script/ir_builder/ir/ir.cc
index d1840e2fef..55e4f1b8f4 100644
--- a/src/script/ir_builder/ir/ir.cc
+++ b/src/script/ir_builder/ir/ir.cc
@@ -34,7 +34,7 @@ using tvm::script::ir_builder::details::Namer;
TVM_STATIC_IR_FUNCTOR(Namer, vtable)
.set_dispatch<tvm::VarNode>([](const ffi::ObjectRef& node, ffi::String
name) -> void {
VarNode* var = const_cast<VarNode*>(node.as<VarNode>());
- var->name_hint = name;
+ var->name = name;
});
IRModuleFrame IRModule() {
diff --git a/src/target/llvm/codegen_cpu.cc b/src/target/llvm/codegen_cpu.cc
index 00ef46c952..f6dcb29cec 100644
--- a/src/target/llvm/codegen_cpu.cc
+++ b/src/target/llvm/codegen_cpu.cc
@@ -546,7 +546,7 @@ void CodeGenCPU::CreateComputeScope(const AttrStmtNode* op)
{
std::vector<llvm::Type*> arg_types;
for (Var v : vargs) {
llvm::Value* value = MakeValue(v);
- value->setName(v->name_hint.c_str());
+ value->setName(v->name.c_str());
arg_values.push_back(value);
arg_types.push_back(value->getType());
}
@@ -561,7 +561,7 @@ void CodeGenCPU::CreateComputeScope(const AttrStmtNode* op)
{
SetTargetAttributes(fcompute);
for (auto it = fcompute->arg_begin(); it != fcompute->arg_end(); it++) {
const Var& var = vargs[std::distance(fcompute->arg_begin(), it)];
- it->setName(std::string(var->name_hint));
+ it->setName(std::string(var->name));
}
llvm::BasicBlock* compute_call_end =
CheckCallSuccess(builder_->CreateCall(fcompute, arg_values));
@@ -635,8 +635,7 @@ void CodeGenCPU::UnpackClosureData(TypedPointer cdata,
const ffi::Array<Var>& vf
llvm::Type* field_type = cdata.type->getStructElementType(i);
llvm::Value* field_addr =
builder_->CreateInBoundsGEP(cdata.type, cdata.addr, {ConstInt32(0),
ConstInt32(i)});
- llvm::Value* load =
- builder_->CreateLoad(field_type, field_addr,
std::string(vfields[i]->name_hint));
+ llvm::Value* load = builder_->CreateLoad(field_type, field_addr,
std::string(vfields[i]->name));
(*vmap)[vfields[i].get()] = load;
}
}
@@ -1190,7 +1189,7 @@ void CodeGenCPU::VisitStmt_(const ForNode* op) {
if (parallel_env_.penv == nullptr) {
auto copy_node = For(ffi::make_object<ForNode>(*op));
CreateParallelLaunch(copy_node, 0,
- std::string("loop_parallel_") +
op->loop_var->name_hint.c_str());
+ std::string("loop_parallel_") +
op->loop_var->name.c_str());
} else {
// already in parallel env.
TVM_FFI_ICHECK(parallel_env_.task_id.defined());
diff --git a/src/target/llvm/codegen_llvm.cc b/src/target/llvm/codegen_llvm.cc
index 89aeb9bde9..4e8c3beda5 100644
--- a/src/target/llvm/codegen_llvm.cc
+++ b/src/target/llvm/codegen_llvm.cc
@@ -324,7 +324,7 @@ void CodeGenLLVM::AddFunctionInternal(const GlobalVar&
gvar, const PrimFunc& f)
llvm::Argument* v = &(*arg_it);
const Var& var = f->params[i];
var_map_[var.get()] = v;
- v->setName(std::string(var->name_hint));
+ v->setName(std::string(var->name));
if (is_restricted_) {
if (var->ty.as<PointerTypeNode>() && !alias_var_set_.count(var.get())) {
// set non alias.
@@ -875,7 +875,7 @@ llvm::Value*
CodeGenLLVM::CreateVecConcat(std::vector<llvm::Value*> vecs) {
void CodeGenLLVM::CreateSerialFor(llvm::Value* begin, llvm::Value* end,
llvm::Value* stride,
const PrimVar& loop_var, const Stmt& body) {
llvm::BasicBlock* pre_block = builder_->GetInsertBlock();
- std::string loop_var_name = loop_var->name_hint;
+ std::string loop_var_name = loop_var->name;
llvm::LLVMContext* ctx = llvm_target_->GetContext();
auto* for_begin = llvm::BasicBlock::Create(*ctx, "for_begin_" +
loop_var_name, function_);
auto* for_body = llvm::BasicBlock::Create(*ctx, "for_body_" + loop_var_name,
function_);
@@ -1021,7 +1021,7 @@ CodeGenLLVM::TypedPointer
CodeGenLLVM::CreateBufferPtr(llvm::Value* buffer_ptr,
llvm::Value* CodeGenLLVM::GetVarValue(const VarNode* v) const {
auto it = var_map_.find(v);
- TVM_FFI_ICHECK(it != var_map_.end()) << "cannot find variable " <<
v->name_hint;
+ TVM_FFI_ICHECK(it != var_map_.end()) << "cannot find variable " << v->name;
return it->second;
}
@@ -2191,7 +2191,7 @@ void CodeGenLLVM::VisitStmt_(const BindNode* op) {
<< ", but is being bound to expression with type " << op->value->ty;
auto* llvm_type = GetLLVMType(v->ty);
if (llvm_type != value->getType()) {
- value->setName((v->name_hint + "_void_ptr").c_str());
+ value->setName((v->name + "_void_ptr").c_str());
value = builder_->CreatePointerCast(value, llvm_type);
}
}
@@ -2300,7 +2300,7 @@ void CodeGenLLVM::AddDebugInformation(llvm::Function*
f_llvm,
void CodeGenLLVM::AddDebugInformation(llvm::Value* llvm_value, const Var&
tir_var,
llvm::Instruction* insert_before) {
- llvm_value->setName(tir_var->name_hint.c_str());
+ llvm_value->setName(tir_var->name.c_str());
if (!di_subprogram_) return;
@@ -2308,7 +2308,7 @@ void CodeGenLLVM::AddDebugInformation(llvm::Value*
llvm_value, const Var& tir_va
// no invalid dtypes
if (!dbg_dtype) return;
auto local_var = dbg_info_->di_builder_->createAutoVariable(
- di_subprogram_, std::string(tir_var->name_hint), dbg_info_->file_, 0,
dbg_dtype);
+ di_subprogram_, std::string(tir_var->name), dbg_info_->file_, 0,
dbg_dtype);
auto* di_loc = llvm::DILocation::get(*llvm_target_->GetContext(), 0, 0,
di_subprogram_);
diff --git a/src/target/source/codegen_source_base.cc
b/src/target/source/codegen_source_base.cc
index f9635b047a..9efb3815a3 100644
--- a/src/target/source/codegen_source_base.cc
+++ b/src/target/source/codegen_source_base.cc
@@ -55,8 +55,8 @@ std::string CodeGenSourceBase::SSAGetID(std::string src,
const Type& t) {
}
std::string CodeGenSourceBase::AllocVarID(const tirx::VarNode* v) {
- TVM_FFI_ICHECK(!var_idmap_.count(v)) << "Need input to be in SSA form dup "
<< v->name_hint;
- std::string key = v->name_hint;
+ TVM_FFI_ICHECK(!var_idmap_.count(v)) << "Need input to be in SSA form dup "
<< v->name;
+ std::string key = v->name;
std::string vid = name_supply_->FreshName(key);
std::replace(vid.begin(), vid.end(), ':', '_');
std::replace(vid.begin(), vid.end(), '-', '_');
@@ -67,7 +67,7 @@ std::string CodeGenSourceBase::AllocVarID(const
tirx::VarNode* v) {
std::string CodeGenSourceBase::GetVarID(const tirx::VarNode* v) const {
auto it = var_idmap_.find(v);
- TVM_FFI_ICHECK(it != var_idmap_.end()) << "Find undefined Variable " <<
v->name_hint;
+ TVM_FFI_ICHECK(it != var_idmap_.end()) << "Find undefined Variable " <<
v->name;
return it->second;
}
diff --git a/src/te/operation/create_primfunc.cc
b/src/te/operation/create_primfunc.cc
index 926676fbe5..ec485da363 100644
--- a/src/te/operation/create_primfunc.cc
+++ b/src/te/operation/create_primfunc.cc
@@ -499,8 +499,8 @@ Stmt GenerateStmtFromCompute(const te::ComputeOp&
compute_op, CreateFuncInfo* in
bool first_times_define =
std::find(axes_levels[i].begin(), axes_levels[i].end(), axis) !=
axes_levels[i].end();
if (first_times_define) {
- Var loop_var = Var(axis->var->name_hint, index_type);
- Var block_var("v_" + axis->var->name_hint, index_type);
+ Var loop_var = Var(axis->var->name, index_type);
+ Var block_var("v_" + axis->var->name, index_type);
PrimExpr min = axis->dom->min;
PrimExpr extent = axis->dom->extent;
if (i > 0) {
@@ -518,7 +518,7 @@ Stmt GenerateStmtFromCompute(const te::ComputeOp&
compute_op, CreateFuncInfo* in
TVM_FFI_ICHECK_GT(i, 0);
TVM_FFI_ICHECK(scopes[i - 1].axes_remap.count(axis->var));
PrimExpr prev_binding = scopes[i - 1].axes_remap.at(axis->var);
- Var block_var("v_" + axis->var->name_hint, index_type);
+ Var block_var("v_" + axis->var->name, index_type);
Range dom = Range::FromMinExtent(prev_binding, MakeConst(index_type,
1));
IterVar new_block_iter(dom, block_var.as_or_throw<PrimVar>(),
axis->iter_type,
axis->thread_tag, axis->span);
diff --git a/src/tirx/analysis/var_use_def_analysis.cc
b/src/tirx/analysis/var_use_def_analysis.cc
index ab8c66dba3..12e5e451a6 100644
--- a/src/tirx/analysis/var_use_def_analysis.cc
+++ b/src/tirx/analysis/var_use_def_analysis.cc
@@ -150,9 +150,9 @@ void VarUseDefAnalyzer::VisitBuffer(const Buffer& buffer) {
void VarUseDefAnalyzer::HandleDef(const Var& var) {
auto v = var.get();
TVM_FFI_ICHECK(!def_count_.count(v))
- << "variable " << v->name_hint << " has already been defined, the Stmt
is not SSA";
+ << "variable " << v->name << " has already been defined, the Stmt is not
SSA";
TVM_FFI_ICHECK(!use_count_.count(v))
- << "variable " << v->name_hint << " has been used before definition!";
+ << "variable " << v->name << " has been used before definition!";
use_count_[v] = 0;
def_count_[v] = 1;
}
diff --git a/src/tirx/analysis/verify_well_formed.cc
b/src/tirx/analysis/verify_well_formed.cc
index 8603d01eb8..59e79d6f4b 100644
--- a/src/tirx/analysis/verify_well_formed.cc
+++ b/src/tirx/analysis/verify_well_formed.cc
@@ -77,13 +77,13 @@ class BlockVarAccessVerifier : public StmtExprVisitor {
if (it->second == 0) {
TVM_FFI_THROW(InternalError)
<< "Well-formedness check failed: "
- << "Loop iterator var " << op->name_hint << " is defined outside
of any block, "
+ << "Loop iterator var " << op->name << " is defined outside of
any block, "
<< "but is used inside the non-opaque current block \""
<< block_stack_.back()->name_hint << "\".";
} else {
TVM_FFI_THROW(InternalError)
<< "Well-formedness check failed: "
- << "Loop iterator var " << op->name_hint << " is defined in
block \""
+ << "Loop iterator var " << op->name << " is defined in block \""
<< block_stack_[it->second - 1]->name_hint << "\", "
<< "but is used inside the non-opaque current block \""
<< block_stack_.back()->name_hint << "\".";
@@ -184,7 +184,7 @@ class UndefinedVarVerifier : public
Verifier<UndefinedVarVerifier> {
auto verify = Verify(it == currently_defined_.end() ||
redefine_is_allowed);
verify << "ValueError: "
<< "TIR is ill-formed, "
- << "due to multiple nested definitions of variable " <<
var->name_hint << ".";
+ << "due to multiple nested definitions of variable " << var->name
<< ".";
if (it != currently_defined_.end()) {
verify << " It was first defined at " << it->second << ", and was
re-defined at " << path;
}
@@ -195,7 +195,7 @@ class UndefinedVarVerifier : public
Verifier<UndefinedVarVerifier> {
auto verify = Verify(it == previously_defined_.end() ||
redefine_is_allowed);
verify << "ValueError: "
<< "TIR is ill-formed, "
- << "due to multiple definitions of variable " << var->name_hint
<< ".";
+ << "due to multiple definitions of variable " << var->name << ".";
if (it != previously_defined_.end()) {
verify << " It was first defined at " << it->second << ", and was
later re-defined at "
<< path;
@@ -218,7 +218,7 @@ class UndefinedVarVerifier : public
Verifier<UndefinedVarVerifier> {
auto active_def = currently_defined_.find(var);
auto verify = Verify(active_def != currently_defined_.end());
verify << "ValueError: "
- << "Invalid use of undefined variable " << var->name_hint << " at "
<< path << ".";
+ << "Invalid use of undefined variable " << var->name << " at " <<
path << ".";
// Check if there was a previous definition, and append the
// location to the error message if there was. This is to aid in
@@ -339,8 +339,8 @@ class SingleEnvThreadVerifier : public
Verifier<SingleEnvThreadVerifier> {
<< "While multiple tirx::AttrStmt may define the same environment
thread, "
<< "all definitions within a single PrimFunc must share the same
tirx::Var. "
<< "Binding of environment thread \"" << iter_var->thread_tag
- << "\" to the TIR variable " << iter_var->var->name_hint << " at "
<< path
- << " conflicts with the previous binding to the TIR variable " <<
prev_var->name_hint
+ << "\" to the TIR variable " << iter_var->var->name << " at " <<
path
+ << " conflicts with the previous binding to the TIR variable " <<
prev_var->name
<< " at " << path;
} else {
env_thread_vars_.insert({iter_var->thread_tag, {iter_var->var, path}});
diff --git a/src/tirx/ir/buffer.cc b/src/tirx/ir/buffer.cc
index 6374c53aac..3d52c94d86 100644
--- a/src/tirx/ir/buffer.cc
+++ b/src/tirx/ir/buffer.cc
@@ -604,11 +604,11 @@ Buffer::Buffer(Var data, PrimType dtype,
ffi::Array<PrimExpr> shape, ffi::Array<
// pointer. Should be done alongside extensions to StmtExprMutator
// to more easily handle buffer/buffer_var updates.
TVM_FFI_ICHECK(!data->ty.IsMissing())
- << "Variable " << data->name_hint << " is missing a type annotation.";
+ << "Variable " << data->name << " is missing a type annotation.";
TVM_FFI_ICHECK(data->ty.as<PointerTypeNode>())
- << "Variable " << data->name_hint << " is not a pointer.";
+ << "Variable " << data->name << " is not a pointer.";
TVM_FFI_ICHECK(data->ty.as<PointerTypeNode>()->element_type.as<PrimTypeNode>())
- << "Variable " << data->name_hint << " does not point to a primitive.";
+ << "Variable " << data->name << " does not point to a primitive.";
ValidateAxisSeparators(axis_separators, shape.size());
diff --git a/src/tirx/ir/index_map.cc b/src/tirx/ir/index_map.cc
index 5ddd6defa9..e4e473990d 100644
--- a/src/tirx/ir/index_map.cc
+++ b/src/tirx/ir/index_map.cc
@@ -87,7 +87,7 @@ std::pair<IndexMap, PrimExpr> IndexMapInverseImpl(const
IndexMap& self,
// input index as (X.outer, X.inner).
std::string name;
if (auto var = index.as<PrimVar>()) {
- name = var.value()->name_hint;
+ name = var.value()->name;
} else {
name = "axis" + std::to_string(i);
}
@@ -387,9 +387,8 @@ IndexMap IndexMap::RenameVariables(
// The name of the variable is pre-defined.
continue;
}
- ffi::String unique_name =
- name_supply->FreshName(initial_index->name_hint, /*add_prefix=*/false);
- if (unique_name != initial_index->name_hint) {
+ ffi::String unique_name = name_supply->FreshName(initial_index->name,
/*add_prefix=*/false);
+ if (unique_name != initial_index->name) {
var_remap.Set(initial_index, PrimVar(unique_name));
}
}
@@ -419,7 +418,7 @@ std::string IndexMap2PythonLambdaExpr(const
ffi::Array<PrimVar>& initial_indices
std::ostringstream oss;
auto print_expr = [&oss](const PrimExpr& expr) {
if (auto var = expr.as<PrimVar>()) {
- oss << var.value()->name_hint;
+ oss << var.value()->name;
} else {
oss << expr;
}
@@ -429,7 +428,7 @@ std::string IndexMap2PythonLambdaExpr(const
ffi::Array<PrimVar>& initial_indices
if (i != 0) {
oss << ", ";
}
- oss << initial_indices[i]->name_hint;
+ oss << initial_indices[i]->name;
}
oss << ": (";
for (size_t i = 0; i < final_indices.size(); ++i) {
diff --git a/src/tirx/script/builder/ir.cc b/src/tirx/script/builder/ir.cc
index 911eaf5d71..a3af8da702 100644
--- a/src/tirx/script/builder/ir.cc
+++ b/src/tirx/script/builder/ir.cc
@@ -628,8 +628,8 @@ LaunchThreadFrame LaunchThread(Var var, PrimExpr extent) {
if (ffi::Optional<IterVar> opt_iter_var =
opt_frame.value()->env_threads.Get(var)) {
iter_var = opt_iter_var.value();
} else {
- TVM_FFI_THROW(InternalError) << "ValueError: " << var->name_hint
- << " is not an env_thread created using
T.env_thread.";
+ TVM_FFI_THROW(InternalError)
+ << "ValueError: " << var->name << " is not an env_thread created
using T.env_thread.";
}
} else {
TVM_FFI_THROW(InternalError) << "LaunchThread can only be used inside a
PrimFunc";
diff --git a/src/tirx/script/printer/buffer.cc
b/src/tirx/script/printer/buffer.cc
index 14bcba4d64..3a43f6fdd4 100644
--- a/src/tirx/script/printer/buffer.cc
+++ b/src/tirx/script/printer/buffer.cc
@@ -132,7 +132,7 @@ ffi::Map<ffi::String, ExprDoc> BufferAttrs(tirx::Buffer
buffer, const AccessPath
return d->AsDoc<ExprDoc>(buffer, buffer_p)
->Attr("strides")[{LiteralDoc::Int(i, std::nullopt)}];
})) {
- results.push_back(LiteralDoc::Str(e.as_or_throw<Var>()->name_hint,
e_p));
+ results.push_back(LiteralDoc::Str(e.as_or_throw<Var>()->name, e_p));
continue;
}
}
diff --git a/src/tirx/script/printer/expr.cc b/src/tirx/script/printer/expr.cc
index f5577e8adc..9beb59799e 100644
--- a/src/tirx/script/printer/expr.cc
+++ b/src/tirx/script/printer/expr.cc
@@ -74,14 +74,14 @@ Doc PrintVar(const tirx::Var& var, const AccessPath& var_p,
const IRDocsifier& d
ExprDoc rhs = PrintVarCreation(var, var_p, d);
opt_f.value()->stmts.push_back(AssignDoc(lhs, rhs, std::nullopt));
} else {
- LOG(WARNING) << "Didn't find variable definition for: " <<
var->name_hint;
+ LOG(WARNING) << "Didn't find variable definition for: " << var->name;
}
}
if (ffi::Optional<ExprDoc> doc = d->GetVarDoc(var)) {
return doc.value();
}
TVM_FFI_THROW(InternalError) << "IndexError: Variable is not defined in the
environment: "
- << var->name_hint;
+ << var->name;
TVM_FFI_UNREACHABLE();
}
@@ -93,7 +93,7 @@ TVM_STATIC_IR_FUNCTOR(IRDocsifier, vtable) //
if (!d->IsVarDefined(var)) {
ExprDoc ann = d->AsDoc<ExprDoc>(var->ty, p->Attr("ty"));
Frame f = d->frames.back();
- ExprDoc lhs = d->Define(var, f, var->name_hint.empty() ? "v" :
var->name_hint);
+ ExprDoc lhs = d->Define(var, f, var->name.empty() ? "v" : var->name);
f->stmts.push_back(AssignDoc(lhs, std::nullopt, ann));
}
return d->GetVarDoc(var).value();
diff --git a/src/tirx/script/printer/stmt.cc b/src/tirx/script/printer/stmt.cc
index cf49aa4972..22e86d44e6 100644
--- a/src/tirx/script/printer/stmt.cc
+++ b/src/tirx/script/printer/stmt.cc
@@ -208,7 +208,7 @@ TVM_STATIC_IR_FUNCTOR(IRDocsifier, vtable)
.set_dispatch<tirx::Bind>("", [](tirx::Bind stmt, AccessPath p,
IRDocsifier d) -> Doc {
// Step 1. Type annotation
TVM_FFI_ICHECK(!stmt->var->ty.IsMissing())
- << "Type annotation is required for variable: " <<
stmt->var->name_hint;
+ << "Type annotation is required for variable: " << stmt->var->name;
ffi::Optional<ExprDoc> type_doc = d->AsDoc<ExprDoc>(stmt->var->ty, //
p->Attr("var")->Attr("ty"));
if (const auto* tuple_type = stmt->var->ty.as<TupleTypeNode>()) {
diff --git a/src/tirx/script/printer/utils.h b/src/tirx/script/printer/utils.h
index 940577d666..b6e07ffd67 100644
--- a/src/tirx/script/printer/utils.h
+++ b/src/tirx/script/printer/utils.h
@@ -91,7 +91,7 @@ inline ExprDoc DefineVar(const tirx::Var& var, const Frame&
frame, const IRDocsi
if (ffi::Optional<ExprDoc> doc = d->GetVarDoc(var)) {
return doc.value();
}
- return d->Define(var, frame, var->name_hint.empty() ? "v" : var->name_hint);
+ return d->Define(var, frame, var->name.empty() ? "v" : var->name);
}
/*!
diff --git a/src/tirx/transform/ir_utils.cc b/src/tirx/transform/ir_utils.cc
index 020b5913e6..5ced08b335 100644
--- a/src/tirx/transform/ir_utils.cc
+++ b/src/tirx/transform/ir_utils.cc
@@ -455,7 +455,7 @@ class IRConvertSSA final : public StmtExprMutator {
it != function_scope_var_remap_.end()) {
var = it->second;
} else if (defined_.count(var.get())) {
- Var new_var(var->name_hint, var->ty);
+ Var new_var(var->name, var->ty);
function_scope_var_remap_.insert({var.get(), new_var});
var = new_var;
@@ -527,7 +527,7 @@ class IRConvertSSA final : public StmtExprMutator {
};
/*! \brief Create a new variable with the same name and type as the
original. */
- static Var MakeNewVar(const Var& old_var) { return Var(old_var->name_hint,
old_var->ty); }
+ static Var MakeNewVar(const Var& old_var) { return Var(old_var->name,
old_var->ty); }
/*! \brief Push a variable remap to the current scope and the var_remap_
stack. */
void PushVarRemap(const Var& old_var, const Var& new_var) {
diff --git a/src/tirx/transform/ir_utils.h b/src/tirx/transform/ir_utils.h
index ed8a95ca0e..1313fe8918 100644
--- a/src/tirx/transform/ir_utils.h
+++ b/src/tirx/transform/ir_utils.h
@@ -114,7 +114,7 @@ inline PrimExpr TVMStructGet(PrimType type, Var handle, int
index,
inline Call AddressOffset(Var handle, PrimType dtype, int offset) {
PrimExpr offset_expr = IntImm::Int32(offset * dtype.lanes());
ffi::Array<PrimExpr> shape = {offset_expr + 1};
- Buffer dummy_buf(handle, dtype, shape, {}, 0, handle->name_hint, 0, 0,
kDefault, {}, Span(),
+ Buffer dummy_buf(handle, dtype, shape, {}, 0, handle->name, 0, 0, kDefault,
{}, Span(),
std::nullopt);
BufferLoad buf_load(dummy_buf, {offset_expr});
@@ -135,7 +135,7 @@ inline Call AddressOffset(Var handle, PrimType dtype,
PrimExpr offset) {
}
ffi::Array<PrimExpr> shape = {offset + 1};
- Buffer dummy_buf(handle, dtype.WithLanes(1), shape, {}, 0,
handle->name_hint, 0, 0, kDefault, {},
+ Buffer dummy_buf(handle, dtype.WithLanes(1), shape, {}, 0, handle->name, 0,
0, kDefault, {},
Span(), std::nullopt);
BufferLoad buf_load(dummy_buf, {offset});
diff --git a/src/tirx/transform/lower_intrin.cc
b/src/tirx/transform/lower_intrin.cc
index 6809024e21..f08bf49ddd 100644
--- a/src/tirx/transform/lower_intrin.cc
+++ b/src/tirx/transform/lower_intrin.cc
@@ -76,7 +76,7 @@ static Expr LowerAccessPtr(const CallNode* call) {
offset = offset * IntImm(offset_ty, dtype.lanes());
offset = Ramp(offset, IntImm(offset_ty, 1), dtype.lanes());
}
- Buffer dummy_buf(buffer_var, dtype.WithLanes(1), {offset + 1}, {}, 0,
buffer_var->name_hint, 0, 0,
+ Buffer dummy_buf(buffer_var, dtype.WithLanes(1), {offset + 1}, {}, 0,
buffer_var->name, 0, 0,
kDefault);
BufferLoad buf_load(dummy_buf, {offset});
return Call(call->ty, builtin::address_of(), {buf_load});
diff --git a/src/tirx/transform/lower_warp_memory.cc
b/src/tirx/transform/lower_warp_memory.cc
index 2e553b26fb..62eb5f4798 100644
--- a/src/tirx/transform/lower_warp_memory.cc
+++ b/src/tirx/transform/lower_warp_memory.cc
@@ -279,7 +279,7 @@ class WarpAccessRewriter : protected StmtExprMutator {
alloc_size = warp_group_ * factor;
Buffer new_buf(op->buffer->data, op->buffer->dtype,
{IntImm::Int32(alloc_size / width_)}, {},
- PrimExpr(), op->buffer->data->name_hint, 0, 0,
BufferType::kDefault);
+ PrimExpr(), op->buffer->data->name, 0, 0,
BufferType::kDefault);
Stmt rewritten_body = this->VisitStmt(body);
return SeqStmt::Flatten(AllocBuffer(new_buf, op->annotations),
rewritten_body);
}
diff --git a/src/tirx/transform/narrow_datatype.cc
b/src/tirx/transform/narrow_datatype.cc
index 4b58c1d838..fdcc2eac67 100644
--- a/src/tirx/transform/narrow_datatype.cc
+++ b/src/tirx/transform/narrow_datatype.cc
@@ -245,7 +245,7 @@ class NarrowDataTypeRewriter : public IndexDataTypeRewriter
{
Expr VisitExpr_(const VarNode* op) final {
if (auto it = visitor_.vmap.find(op); !var_remap_.count(op) && it !=
visitor_.vmap.end()) {
- var_remap_[op] = Var(op->name_hint, it->second);
+ var_remap_[op] = Var(op->name, it->second);
}
return Parent::VisitExpr_(op);
}
diff --git a/src/tirx/transform/split_host_device.cc
b/src/tirx/transform/split_host_device.cc
index a73ed07cb7..a54bc3e0fd 100644
--- a/src/tirx/transform/split_host_device.cc
+++ b/src/tirx/transform/split_host_device.cc
@@ -143,7 +143,7 @@ class HostDeviceSplitter : public StmtMutator {
bool is_handle = var->ty.as<PointerTypeNode>() != nullptr;
return std::tuple{
!is_handle,
- var->name_hint,
+ var->name,
};
};
return sort_key(a) < sort_key(b);
@@ -330,7 +330,7 @@ class DeviceInfoCollector : public StmtVisitor {
thread_tag = iv.value()->thread_tag;
TVM_FFI_ICHECK_NE(thread_tag.length(), 0U);
} else if (auto var = op->node.as<Var>()) {
- thread_tag = var.value()->name_hint;
+ thread_tag = var.value()->name;
} else {
TVM_FFI_THROW(TypeError) << "thread_extent node must be an IterVar or
Var, but was "
<< op->node.GetTypeKey();
diff --git a/src/tirx/transform/storage_rewrite.cc
b/src/tirx/transform/storage_rewrite.cc
index ef17566f37..0b0697446b 100644
--- a/src/tirx/transform/storage_rewrite.cc
+++ b/src/tirx/transform/storage_rewrite.cc
@@ -183,7 +183,7 @@ class LinearAccessPatternFinder final : public
StmtExprVisitor {
// Directly reference to the variable count as a read.
auto it = alloc_info_.find(buf);
if (it != alloc_info_.end() && it->second.alloc) {
- TVM_FFI_ICHECK_LT(it->second.level, scope_.size()) << " buf=" <<
buf->name_hint;
+ TVM_FFI_ICHECK_LT(it->second.level, scope_.size()) << " buf=" <<
buf->name;
scope_[it->second.level].touched.push_back(buf);
}
}
@@ -377,7 +377,7 @@ class InplaceOpVerifier : public StmtExprVisitor {
return;
}
TVM_FFI_ICHECK_EQ(store_->indices.size(), op->indices.size())
- << "Store/Load occur to the same buffer " << buf->name_hint
+ << "Store/Load occur to the same buffer " << buf->name
<< " with differing number of indices";
for (size_t i = 0; i < store_->indices.size(); i++) {
if (!tirx::ExprDeepEqual()(store_->indices[i], op->indices[i])) {
@@ -461,15 +461,15 @@ class StoragePlanRewriter : public StmtExprMutator {
if (it != buffer_remap_.end()) {
TVM_FFI_ICHECK_EQ(it->second->data.get(), new_backing_array.get())
<< "Cannot remap buffer " << buf->name << " to use backing array "
- << new_backing_array->name_hint << ", previously used backing array "
- << it->second->data->name_hint;
+ << new_backing_array->name << ", previously used backing array "
+ << it->second->data->name;
return it->second;
}
Buffer remapped =
Buffer(new_backing_array, buf->dtype, buf->shape, buf->strides,
buf->elem_offset,
- new_backing_array->name_hint, buf->data_alignment,
buf->offset_factor,
- buf->buffer_type, buf->axis_separators, buf->span, buf->layout,
buf->allocated_addr);
+ new_backing_array->name, buf->data_alignment,
buf->offset_factor, buf->buffer_type,
+ buf->axis_separators, buf->span, buf->layout,
buf->allocated_addr);
buffer_remap_[key] = remapped;
return remapped;
}
@@ -741,7 +741,7 @@ class StoragePlanRewriter : public StmtExprMutator {
PrimExpr combo_size;
for (const AllocBufferNode* op : e->allocs) {
TVM_FFI_ICHECK_EQ(op->buffer->shape.size(), 1)
- << "Buffer var " << op->buffer->data->name_hint
+ << "Buffer var " << op->buffer->data->name
<< " was identified as a re-usable allocation, but has " <<
op->buffer->shape.size()
<< " physical dimensions. "
<< "Currently, only flat 1-d memory spaces should be
identified as re-usable "
@@ -774,8 +774,8 @@ class StoragePlanRewriter : public StmtExprMutator {
combo_size = combo_size + IntImm::Int32(1);
}
combo_size = analyzer_->Simplify(combo_size);
- Buffer buf(e->alloc_var, alloc_type, {combo_size}, {}, PrimExpr(),
- e->alloc_var->name_hint, 0, 0, BufferType::kDefault);
+ Buffer buf(e->alloc_var, alloc_type, {combo_size}, {}, PrimExpr(),
e->alloc_var->name, 0,
+ 0, BufferType::kDefault);
ffi::Map<ffi::String, ffi::Any> annotations;
if (e->is_volatile) {
annotations.Set(attr::kVolatile, true);
@@ -812,8 +812,8 @@ class StoragePlanRewriter : public StmtExprMutator {
uint64_t type_bits = e->elem_type.bits() * e->elem_type.lanes();
PrimExpr alloc_size =
MakeConst(e->allocs[0]->buffer->shape[0].ty(), (total_bits + type_bits
- 1) / type_bits);
- Buffer buf(e->alloc_var, e->elem_type, {alloc_size}, {}, PrimExpr(),
e->alloc_var->name_hint, 0,
- 0, BufferType::kDefault);
+ Buffer buf(e->alloc_var, e->elem_type, {alloc_size}, {}, PrimExpr(),
e->alloc_var->name, 0, 0,
+ BufferType::kDefault);
bool any_volatile = e->is_volatile;
for (StorageEntry* child : e->merged_children) {
if (child->is_volatile) any_volatile = true;
@@ -1327,7 +1327,7 @@ class VectorTypeAccessChecker : public StmtExprVisitor {
void OnArrayDeclaration(Var buffer, PrimType element_dtype, PrimExpr extent,
BufferVarInfo::DeclarationLocation
declaration_location) {
TVM_FFI_ICHECK(info_map_.find(buffer.get()) == info_map_.end())
- << "Array declaration of " << buffer->name_hint << " occurred multiple
times.";
+ << "Array declaration of " << buffer->name << " occurred multiple
times.";
if (element_dtype.MatchesCode(DLDataTypeCode::kDLBool)) {
element_dtype = PrimType::Int(8, element_dtype.lanes());
@@ -1350,7 +1350,7 @@ class VectorTypeAccessChecker : public StmtExprVisitor {
void OnArrayAccess(PrimType value_dtype, const VarNode* buffer,
const ffi::Array<PrimExpr>& indices, bool is_buffer_load)
{
auto it = info_map_.find(buffer);
- TVM_FFI_ICHECK(it != info_map_.end()) << "Load/Store of buffer " <<
buffer->name_hint << " ("
+ TVM_FFI_ICHECK(it != info_map_.end()) << "Load/Store of buffer " <<
buffer->name << " ("
<< buffer << ") occurred before its
declaration.";
if (value_dtype.IsScalableVector()) {
@@ -1367,8 +1367,7 @@ class VectorTypeAccessChecker : public StmtExprVisitor {
if (var_info.element_dtype.IsVoid()) {
TVM_FFI_ICHECK(allow_untyped_pointers_)
- << "Variable " << buffer->name_hint
- << " was missing a type annotation in its declaration";
+ << "Variable " << buffer->name << " was missing a type annotation in
its declaration";
var_info.element_dtype = value_dtype.WithLanes(1);
}
@@ -1516,7 +1515,7 @@ class VectorTypeRewriter : public StmtExprMutator {
PrimType preferred = var_info.get_preferred_dtype();
if (preferred != var_info.element_dtype && (rewrite_mask &
var_info.declaration_location)) {
Var old_buffer_var = var_info.var;
- Var new_buffer_var(old_buffer_var->name_hint,
+ Var new_buffer_var(old_buffer_var->name,
PointerType(preferred,
GetPtrStorageScope(old_buffer_var)),
old_buffer_var->span);
diff --git a/src/tirx/transform/tile_primitive_dispatch.cc
b/src/tirx/transform/tile_primitive_dispatch.cc
index 06c522c254..de80b9cfd3 100644
--- a/src/tirx/transform/tile_primitive_dispatch.cc
+++ b/src/tirx/transform/tile_primitive_dispatch.cc
@@ -643,7 +643,7 @@ class TilePrimitiveDispatcher : public StmtExprMutator {
ScopeIdDefVerifier verifier;
TVM_FFI_ICHECK(verifier.Verify(defs)) << "Inconsistent ScopeIdDef";
- auto is_implicit = [](const Var& v) { return v->name_hint.empty(); };
+ auto is_implicit = [](const Var& v) { return v->name.empty(); };
for (const auto& g : gathered) {
ScopeIdDef def = g.def;
// Deferred extents: resolved via closure into verifier.id_set.
diff --git a/src/tirx/transform/tvm_ffi_binder.cc
b/src/tirx/transform/tvm_ffi_binder.cc
index 26fe4a193d..8c487c842f 100644
--- a/src/tirx/transform/tvm_ffi_binder.cc
+++ b/src/tirx/transform/tvm_ffi_binder.cc
@@ -67,7 +67,7 @@ TVMFFIABIBuilder::TVMFFIABIBuilder(const ffi::String&
func_name, const ffi::Arra
if (j > 0) os << ", ";
std::ostringstream shape_os;
if (auto var = buf->shape[j].as<PrimVar>()) {
- shape_os << ((*var)->name_hint.empty() ? "v" :
(*var)->name_hint.c_str());
+ shape_os << ((*var)->name.empty() ? "v" : (*var)->name.c_str());
} else {
shape_os << buf->shape[j];
}
@@ -76,13 +76,13 @@ TVMFFIABIBuilder::TVMFFIABIBuilder(const ffi::String&
func_name, const ffi::Arra
os << "], " << buf->dtype->dtype << ")";
param_names_[static_cast<int>(i)] = buf_name;
} else {
- os << param->name_hint << ": ";
+ os << param->name << ": ";
if (const auto* prim_type = param->ty.as<PrimTypeNode>()) {
os << prim_type->dtype;
} else {
os << param->ty;
}
- param_names_[static_cast<int>(i)] = param->name_hint;
+ param_names_[static_cast<int>(i)] = param->name;
}
}
os << ")";
@@ -275,7 +275,7 @@ bool TVMFFIABIBuilder::BindPointer(const Var& arg, const
Expr& value,
* names.
*
* Uses ExprFunctor for generic dispatch over all expression types.
- * The default TIR printer sanitizes Var name_hints (e.g. "B.shape[0]" ->
"B_shape_0_")
+ * The default TIR printer sanitizes Var names (e.g. "B.shape[0]" ->
"B_shape_0_")
* and adds type annotations (e.g. T.int64(1)). This functor preserves
original path
* names and uses plain integer formatting for human-readable error messages.
*/
@@ -332,7 +332,7 @@ void TVMFFIABIBuilder::RenderPendingAsserts() {
ffi::String path = RenderAccessPath(it->second.first_def_path);
if (!path.empty()) return std::string(path);
}
- return std::string(v->name_hint);
+ return std::string(v->name);
});
for (auto& pending : pending_const_asserts_) {
@@ -549,7 +549,7 @@ void TVMFFIABIBuilder::DecodeParam(int param_index) {
Var param = params_[param_index];
// Extract type_index from packed_args
- Var type_index(param->name_hint + ".type_index", PrimType::Int(32));
+ Var type_index(param->name + ".type_index", PrimType::Int(32));
init_nest_.push_back(Bind(type_index, Call(PrimType::Int(32),
builtin::tvm_struct_get(),
{v_packed_args_,
IntImm::Int32(param_index),
IntImm::Int32(builtin::kTVMFFIAnyTypeIndex)})
@@ -603,8 +603,8 @@ void TVMFFIABIBuilder::DecodeAllParams() {
ffi::reflection::AccessPath param_path =
ffi::reflection::AccessPath::Root()
->Extend(AccessStep::ArrayItem(i))
->Attr(ffi::String(buffer->name));
- DecodeParamDLTensor(buffer, device_type_, device_id_, param,
- func_name_ + "." + param->name_hint, param_path);
+ DecodeParamDLTensor(buffer, device_type_, device_id_, param, func_name_
+ "." + param->name,
+ param_path);
decl_buffers_.push_back(DeclBuffer(buffer));
}
}
diff --git a/src/tirx/transform/unsupported_dtype_legalize.cc
b/src/tirx/transform/unsupported_dtype_legalize.cc
index 2d2e8ed74a..8da543e8f9 100644
--- a/src/tirx/transform/unsupported_dtype_legalize.cc
+++ b/src/tirx/transform/unsupported_dtype_legalize.cc
@@ -119,7 +119,7 @@ class ComputeLegalizePlanner : public StmtExprVisitor {
if (auto* ptr_type = op->buffer->data->ty.as<PointerTypeNode>()) {
storage_scope = ptr_type->storage_scope;
}
- Var buffer_var = Var(op->buffer->data->name_hint, PointerType(dtype,
storage_scope));
+ Var buffer_var = Var(op->buffer->data->name, PointerType(dtype,
storage_scope));
(*var_remap_)[op->buffer->data] = buffer_var;
}
return StmtExprVisitor::VisitStmt_(op);
@@ -583,7 +583,7 @@ class StorageLegalizer : public StmtExprMutator {
if (auto* ptr_type = buf->data->ty.as<PointerTypeNode>()) {
storage_scope = ptr_type->storage_scope;
}
- Var new_data = Var(buf->data->name_hint, PointerType(new_dtype,
storage_scope));
+ Var new_data = Var(buf->data->name, PointerType(new_dtype,
storage_scope));
var_remap_[buf->data] = new_data;
buf = Buffer(new_data, new_dtype, buf->shape, buf->strides,
buf->elem_offset, buf->name,
buf->data_alignment, buf->offset_factor, buf->buffer_type,
buf->axis_separators,
@@ -747,8 +747,8 @@ class StorageLegalizer : public StmtExprMutator {
if (auto* elem_type = ptr_type->element_type.as<PrimTypeNode>()) {
PrimType elem_prim_type = ffi::GetRef<PrimType>(elem_type);
if (MatchType(elem_prim_type)) {
- Var new_var = Var(var->name_hint,
PointerType(GetStorageUIntDType(elem_prim_type),
-
ptr_type->storage_scope));
+ Var new_var = Var(
+ var->name, PointerType(GetStorageUIntDType(elem_prim_type),
ptr_type->storage_scope));
var_remap_[var] = new_var;
return new_var;
}
diff --git a/src/tirx/transform/update_pointer_storage_scope.cc
b/src/tirx/transform/update_pointer_storage_scope.cc
index e19df95561..0c796ff364 100644
--- a/src/tirx/transform/update_pointer_storage_scope.cc
+++ b/src/tirx/transform/update_pointer_storage_scope.cc
@@ -41,7 +41,7 @@ namespace tirx {
Var WithStorageScope(const VarNode* buffer_var, ffi::String storage_scope) {
auto* ptr_type = buffer_var->ty.as<PointerTypeNode>();
TVM_FFI_ICHECK(ptr_type) << "The provided variable is not of pointer type";
- return Var(buffer_var->name_hint, PointerType(ptr_type->element_type,
storage_scope),
+ return Var(buffer_var->name, PointerType(ptr_type->element_type,
storage_scope),
buffer_var->span);
}
diff --git a/src/tirx/transform/vectorize_loop.cc
b/src/tirx/transform/vectorize_loop.cc
index 996bac3464..9872a8c462 100644
--- a/src/tirx/transform/vectorize_loop.cc
+++ b/src/tirx/transform/vectorize_loop.cc
@@ -748,7 +748,7 @@ class Vectorizer : public StmtMutator, public
ExprFunctor<Expr(const Expr&)> {
<< "Let cannot bind the same var to two different values";
}
if (GetLanesOrVScaleFactor(value.ty()) !=
GetLanesOrVScaleFactor(op->value.ty())) {
- Var new_var(op->var->name_hint, value.ty());
+ Var new_var(op->var->name, value.ty());
let_binding_[op->var] = new_var.as_or_throw<PrimExpr>();
return Let(new_var, value, this->VisitPrimExpr(op->body));
} else {
@@ -961,7 +961,7 @@ class Vectorizer : public StmtMutator, public
ExprFunctor<Expr(const Expr&)> {
let_binding_[op->var] = value;
if (GetLanesOrVScaleFactor(value.ty()) !=
GetLanesOrVScaleFactor(prim_value.value().ty())) {
- Var new_var(op->var->name_hint, value.ty());
+ Var new_var(op->var->name, value.ty());
let_binding_[op->var] = new_var.as_or_throw<PrimExpr>();
return Bind(new_var, value);
} else {
@@ -979,7 +979,7 @@ class Vectorizer : public StmtMutator, public
ExprFunctor<Expr(const Expr&)> {
// scalarize the statment
Stmt Scalarize(Stmt stmt) {
PrimType var_ty = var_->ty.as_or_throw<PrimType>();
- Var idx(var_->name_hint + ".s", var_ty);
+ Var idx(var_->name + ".s", var_ty);
stmt = Substitute(stmt, ffi::Map<Var, Expr>{{var_, idx}});
return For(idx.as_or_throw<PrimVar>(), IntImm(var_ty, 0), var_lanes_,
ForKind::kSerial, stmt);
}
@@ -1157,8 +1157,8 @@ class LoopVectorizer : public StmtMutator {
}
PrimExpr num_chunks = ceildiv(fixed_extent, scalable_lanes_index);
- PrimVar outer(op->loop_var->name_hint + ".vla.o", index_dtype);
- PrimVar inner(op->loop_var->name_hint + ".vla.i", lane_dtype);
+ PrimVar outer(op->loop_var->name + ".vla.o", index_dtype);
+ PrimVar inner(op->loop_var->name + ".vla.i", lane_dtype);
PrimExpr inner_index = inner;
if (inner_index.ty() != index_dtype) {
inner_index = Cast(index_dtype, inner_index);
diff --git a/tests/cpp/expr_test.cc b/tests/cpp/expr_test.cc
index 6ad0f81191..1b4dca87e7 100644
--- a/tests/cpp/expr_test.cc
+++ b/tests/cpp/expr_test.cc
@@ -59,14 +59,14 @@ TEST(Expr, VarCopyHelpers) {
Var renamed = var.CopyWithName("y");
EXPECT_FALSE(renamed.same_as(var));
- EXPECT_EQ(renamed->name_hint, "y");
+ EXPECT_EQ(renamed->name, "y");
EXPECT_TRUE(renamed->ty.same_as(pointer_type));
EXPECT_TRUE(renamed->span.same_as(span));
PrimType dtype = PrimType::Int(64);
Var retyped = var.CopyWithDType(dtype);
EXPECT_FALSE(retyped.same_as(var));
- EXPECT_EQ(retyped->name_hint, "x");
+ EXPECT_EQ(retyped->name, "x");
EXPECT_TRUE(retyped->ty.same_as(dtype));
EXPECT_TRUE(retyped->span.same_as(span));
@@ -75,7 +75,7 @@ TEST(Expr, VarCopyHelpers) {
PrimType prim_dtype = PrimType::Float(32);
PrimVar retyped_prim_var = prim_var.CopyWithDType(prim_dtype);
EXPECT_FALSE(retyped_prim_var.same_as(prim_var));
- EXPECT_EQ(retyped_prim_var->name_hint, "i");
+ EXPECT_EQ(retyped_prim_var->name, "i");
EXPECT_TRUE(retyped_prim_var.ty().same_as(prim_dtype));
EXPECT_TRUE(retyped_prim_var->span.same_as(span));
}
diff --git a/tests/python/ir/test_node_reflection.py
b/tests/python/ir/test_node_reflection.py
index d5c260c783..29e222c6db 100644
--- a/tests/python/ir/test_node_reflection.py
+++ b/tests/python/ir/test_node_reflection.py
@@ -80,23 +80,37 @@ _LEGACY_TIRX_VAR_JSON = """{
}"""
+def _make_pre_name_field_json(type_key):
+ graph = json.loads(_LEGACY_TIRX_VAR_JSON)
+ graph["root_index"] = 5
+ var = graph["nodes"][5]
+ var["type"] = type_key
+ var["data"]["name_hint"] = var["data"].pop("name")
+ return json.dumps(graph)
+
+
+_PRE_NAME_FIELD_IR_VAR_JSON = _make_pre_name_field_json("ir.Var")
+_PRE_NAME_FIELD_DATAFLOW_VAR_JSON =
_make_pre_name_field_json("relax.expr.DataflowVar")
+
+
@pytest.mark.parametrize(
- ("legacy_json", "legacy_type", "var_index"),
+ ("legacy_json", "expected_type", "var_index"),
[
- (_LEGACY_RELAX_VAR_JSON, "relax.expr.Var", 7),
- (_LEGACY_TIRX_VAR_JSON, "tirx.Var", 5),
+ (_LEGACY_RELAX_VAR_JSON, "ir.Var", 7),
+ (_LEGACY_TIRX_VAR_JSON, "ir.Var", 5),
+ (_PRE_NAME_FIELD_IR_VAR_JSON, "ir.Var", 5),
+ (_PRE_NAME_FIELD_DATAFLOW_VAR_JSON, "relax.expr.DataflowVar", 5),
],
)
-def test_var_exact_base_legacy_json_graph_rewrite(legacy_json, legacy_type,
var_index):
+def test_var_name_legacy_json_graph_rewrite(legacy_json, expected_type,
var_index):
from tvm.ir.json_compact import upgrade_json
original = json.loads(legacy_json)
expected = copy.deepcopy(original)
- expected["nodes"][var_index]["type"] = "ir.Var"
- if legacy_type == "tirx.Var":
- expected["nodes"][var_index]["data"]["name_hint"] =
expected["nodes"][var_index][
- "data"
- ].pop("name")
+ expected["nodes"][var_index]["type"] = expected_type
+ fields = expected["nodes"][var_index]["data"]
+ if "name_hint" in fields:
+ fields["name"] = fields.pop("name_hint")
upgraded = json.loads(upgrade_json(legacy_json))
assert upgraded == expected
@@ -104,8 +118,10 @@ def
test_var_exact_base_legacy_json_graph_rewrite(legacy_json, legacy_type, var_
assert len(upgraded["nodes"]) == len(original["nodes"])
-def _check_legacy_var(var, source_name, line, end_line, column, end_column):
- assert type(var) is tvm.ir.Var
+def _check_legacy_var(
+ var, source_name, line, end_line, column, end_column,
expected_type=tvm.ir.Var
+):
+ assert type(var) is expected_type
assert var.name == "legacy"
assert var.ty == tvm.ir.PrimType("int64")
assert var.span.source_name.name == source_name
@@ -137,6 +153,22 @@ def test_var_exact_base_legacy_tirx_json_load():
)
[email protected](
+ ("legacy_json", "expected_type"),
+ [
+ (_PRE_NAME_FIELD_IR_VAR_JSON, tvm.ir.Var),
+ (_PRE_NAME_FIELD_DATAFLOW_VAR_JSON, tvm.relax.DataflowVar),
+ ],
+)
+def test_var_name_legacy_json_load(legacy_json, expected_type):
+ restored = tvm.ir.load_json(legacy_json)
+ _check_legacy_var(restored, "legacy_tirx.py", 7, 9, 2, 14, expected_type)
+ graph = json.loads(tvm.ir.save_json(restored))
+ fields = graph["nodes"][graph["root_index"]]["data"]
+ assert "name" in fields
+ assert "name_hint" not in fields
+
+
def test_dataflow_var_json_is_not_migrated_to_canonical_var():
dataflow_var = tvm.relax.DataflowVar("value", tvm.ir.PrimType("int64"))
graph = json.loads(tvm.ir.save_json(dataflow_var))
diff --git a/tests/python/relax/test_expr.py b/tests/python/relax/test_expr.py
index 52b4e9ba16..5a43477702 100644
--- a/tests/python/relax/test_expr.py
+++ b/tests/python/relax/test_expr.py
@@ -60,8 +60,12 @@ def test_var() -> None:
def test_var_name_keyword_compatibility() -> None:
- assert tvm.ir.Var(name="primary").name == "primary"
- assert tvm.ir.Var(name_hint="legacy").name == "legacy"
+ primary = tvm.ir.Var(name="primary")
+ legacy = tvm.ir.Var(name_hint="legacy")
+ assert primary.name == "primary"
+ assert legacy.name == "legacy"
+ assert not hasattr(primary, "name_hint")
+ assert not hasattr(legacy, "name_hint")
with pytest.raises(TypeError, match="Specify either name or name_hint, not
both"):
tvm.ir.Var(name="primary", name_hint="legacy")
@@ -102,8 +106,12 @@ def test_dataflow_var() -> None:
def test_dataflow_var_name_keyword_compatibility() -> None:
- assert rx.DataflowVar(name="primary").name == "primary"
- assert rx.DataflowVar(name_hint="legacy").name == "legacy"
+ primary = rx.DataflowVar(name="primary")
+ legacy = rx.DataflowVar(name_hint="legacy")
+ assert primary.name == "primary"
+ assert legacy.name == "legacy"
+ assert not hasattr(primary, "name_hint")
+ assert not hasattr(legacy, "name_hint")
with pytest.raises(TypeError, match="Specify either name or name_hint, not
both"):
rx.DataflowVar(name="primary", name_hint="legacy")
