[IR] Validate global_symbol on all functions, not just Relax

python/tvm/runtime/_tensor.py

@@ -15,7 +15,7 @@
 # specific language governing permissions and limitations
 # under the License.
 # pylint: disable=invalid-name, unused-import, redefined-outer-name
-# ruff: noqa: E722, F401, RUF005
+# ruff: noqa: F401, RUF005
 """Runtime Tensor API"""
 
 import ctypes

python/tvm/s_tir/dlight/analysis/common_analysis.py

@@ -14,7 +14,6 @@
 # KIND, either express or implied.  See the License for the
 # specific language governing permissions and limitations
 # under the License.
-# ruff: noqa: E722
 
 # pylint: disable=missing-function-docstring, missing-class-docstring
 # pylint: disable=unused-argument, unused-variable

src/relax/analysis/well_formed.cc

@@ -93,12 +93,12 @@ class WellFormedChecker : public relax::ExprVisitor,
         WellFormedChecker(obj.as<IRModule>(), check_struct_info);
 
     if (const auto* mod = obj.as<IRModuleNode>()) {
-      for (const auto& it : mod->functions) {
+      for (const auto& [gvar, base_func] : mod->functions) {
+        well_formed_checker.CheckGlobalVarAndGsymbolConsistency(gvar, base_func);
         // visit relax.Function
-        if (auto* n = it.second.as<FunctionNode>()) {
-          Function func = ffi::GetRef<Function>(n);
-          well_formed_checker.func_name_map_[n] = it.first->name_hint;
-          well_formed_checker.CheckGlobalVarAndGsymbolConsistency(it.first, func);
+        if (auto opt = base_func.as<Function>()) {
+          Function func = opt.value();
+          well_formed_checker.func_name_map_[func.get()] = gvar->name_hint;
           well_formed_checker.VisitExpr(func);
         }
       }
@@ -146,7 +146,7 @@ class WellFormedChecker : public relax::ExprVisitor,
     return "(anonymous function)";
   }
 
-  void CheckGlobalVarAndGsymbolConsistency(GlobalVar var, Function func) {
+  void CheckGlobalVarAndGsymbolConsistency(GlobalVar var, BaseFunc func) {
     // the uniqueness of all global vars are ensured by IRModule->global_var_map_, so do not need
     // to check again
 

tests/python/relax/test_group_gemm_flashinfer.py

@@ -14,7 +14,7 @@
 # KIND, either express or implied.  See the License for the
 # specific language governing permissions and limitations
 # under the License.
-# ruff: noqa: E501, E722, F401, F841, RUF005
+# ruff: noqa: E501, F401, F841, RUF005
 
 """Test for FlashInfer GroupedGemm TVM integration"""
 

tests/python/s_tir/transform/test_s_tir_transform_manifest_shared_memory_local_stage.py

@@ -29,7 +29,7 @@ class MatmulBefore:
     @T.prim_func
     def main(A: T.Buffer((1024, 1024), "float32"), B: T.Buffer((1024, 1024), "float32"), C: T.Buffer((1024, 1024), "float32")) -> None:
         # function attr dict
-        T.func_attr({"global_symbol": "default_function", "tir.noalias": True})
+        T.func_attr({"tir.noalias": True})
         # body
         # with T.sblock("root")
         for blockIdx_y in T.thread_binding(32, thread="blockIdx.y"):
@@ -70,7 +70,7 @@ class MatmulAfter:
     @T.prim_func
     def main(A: T.Buffer((1024, 1024), "float32"), B: T.Buffer((1024, 1024), "float32"), C: T.Buffer((1024, 1024), "float32")) -> None:
         # function attr dict
-        T.func_attr({"global_symbol": "default_function", "tir.noalias": True})
+        T.func_attr({"tir.noalias": True})
         # body
         # with T.sblock("root")
         for blockIdx_y in T.thread_binding(32, thread="blockIdx.y"):

tests/python/tir-base/test_tir_host_func.py

@@ -33,7 +33,6 @@ def main(
     ):
         T.func_attr(
             {
-                "global_symbol": "test",
                 "target": tvm.target.Target("llvm", host="llvm"),
                 "tir.noalias": True,
             }
@@ -59,12 +58,11 @@ def test_host_func():
     func = tvm.te.create_prim_func(
         te_workload.matmul(729, 729, 729, in_dtype="float32", out_dtype="float32")
     )
-    mod = tvm.ir.IRModule({"main": func})
+    mod = tvm.ir.IRModule({"main": func.with_attr("global_symbol", "main")})
     target = tvm.target.Target("cuda")
     mod = tvm.tir.transform.Apply(
         lambda f: f.with_attr(
             {
-                "global_symbol": "test",
                 "tir.is_host_func": True,
             }
         )

tests/python/tir-base/test_tir_intrin.py

@@ -307,7 +307,7 @@ class Module:
     @T.prim_func
     def test_tir_fma(A: T.handle, B: T.handle, C: T.handle, d: T.handle) -> None:
         # function attr dict
-        T.func_attr({"global_symbol": "test_fma", "tir.noalias": True})
+        T.func_attr({"tir.noalias": True})
         n = T.int32()
         stride = T.int32()
         stride_1 = T.int32()

tests/python/tir-transform/test_tir_transform_device_kernel_launch.py

@@ -92,7 +92,7 @@ def main(A: T.Buffer(1, "float32")):
 
         @T.prim_func
         def kernel(A_data: T.handle("float32")):
-            T.func_attr({"target": T.target("cuda"), "global_symbol": "kernel_by_another_name"})
+            T.func_attr({"target": T.target("cuda")})
             A = T.decl_buffer(1, dtype="float32", data=A_data)
             A[0] = 0.0
 
@@ -101,7 +101,7 @@ class Expected:
         @T.prim_func
         def main(A: T.Buffer(1, "float32")):
             T.func_attr({"target": T.target("llvm")})
-            T.call_packed("kernel_by_another_name", A.data)
+            T.call_packed("kernel", A.data)
 
         @T.prim_func
         def kernel(A_data: T.handle("float32")):
@@ -110,7 +110,7 @@ def kernel(A_data: T.handle("float32")):
                     "target": T.target("cuda"),
                     "calling_conv": 2,
                     "tir.kernel_launch_params": [],
-                    "global_symbol": "kernel_by_another_name",
+                    "global_symbol": "kernel",
                     "tir.is_global_func": True,
                 }
             )

tests/scripts/release/make_notes.py

@@ -15,7 +15,6 @@
 # KIND, either express or implied.  See the License for the
 # specific language governing permissions and limitations
 # under the License.
-# ruff: noqa: E722
 
 import argparse
 import csv