[WGSL] Non-entrypoint functions should also be mangled

https://bugs.webkit.org/show_bug.cgi?id=258281
rdar://111004485

Reviewed by Dan Glastonbury.

In the mangling pass, although we visited all functions, we only mangled the
names of the entry points.

* Source/WebGPU/WGSL/MangleNames.cpp:
(WGSL::NameManglerVisitor::run):
* Source/WebGPU/WGSL/tests/valid/global-used-by-callee.wgsl:
* Source/WebGPU/WGSL/tests/valid/name-mangling.wgsl: Added.
* Source/WebGPU/WGSL/tests/valid/var-initialization-with-var.wgsl:

Canonical link: https://commits.webkit.org/266096@main

Source/WebGPU/WGSL/MangleNames.cpp

@@ -103,15 +103,15 @@ class NameManglerVisitor : public AST::Visitor, public ContextProvider<MangledNa
 
 void NameManglerVisitor::run()
 {
-    for (const auto& entrypoint : m_callGraph.entrypoints()) {
-        String originalName = entrypoint.function.name();
-        introduceVariable(entrypoint.function.name(), MangledName::Function);
+    auto& module = m_callGraph.ast();
+    for (auto& function : module.functions()) {
+        String originalName = function.name();
+        introduceVariable(function.name(), MangledName::Function);
         auto it = m_result.entryPoints.find(originalName);
-        RELEASE_ASSERT(it != m_result.entryPoints.end());
-        it->value.mangledName = entrypoint.function.name();
+        if (it != m_result.entryPoints.end())
+            it->value.mangledName = function.name();
     }
 
-    auto& module = m_callGraph.ast();
     for (auto& structure : module.structures())
         visit(structure);
 

Source/WebGPU/WGSL/tests/valid/global-used-by-callee.wgsl

@@ -4,40 +4,40 @@
 var<private> x: i32;
 var<private> y: i32;
 
-// CHECK: int f\(int parameter\d\)
+// CHECK: int function\d\(int parameter\d\)
 fn f() -> i32
 {
     //CHECK: parameter\d
     _ = y;
     return 0;
 }
 
-// CHECK: int g\(\)
+// CHECK: int function\d\(\)
 fn g() -> i32
 {
     let y = 42;
     return 0;
 }
 
-// CHECK: int h\(int parameter\d\)
+// CHECK: int function\d\(int parameter\d\)
 fn h() -> i32
 {
     _ = y;
     return 0;
 }
 
-// CHECK: int i\(int parameter\d\)
+// CHECK: int function\d\(int parameter\d\)
 fn i() -> i32
 {
-    // CHECK: f\(parameter\d\)
+    // CHECK: function\d\(parameter\d\)
     _ = f();
     return 0;
 }
 
-// CHECK: float j\(float parameter\d, int parameter\d\)
+// CHECK: float function\d\(float parameter\d, int parameter\d\)
 fn j(x: f32) -> f32
 {
-    // CHECK: f\(parameter\d\)
+    // CHECK: function\d\(parameter\d\)
     _ = f();
     return x;
 }
@@ -49,18 +49,18 @@ fn main()
     // CHECK: int local\d;
     _ = x;
 
-    // CHECK: f\(local\d\)
+    // CHECK: function\d\(local\d\)
     _ = f();
 
-    // CHECK: g\(\)
+    // CHECK: function\d\(\)
     _ = g();
 
-    // CHECK: h\(local\d\)
+    // CHECK: function\d\(local\d\)
     _ = h();
 
-    // CHECK: i\(local\d\)
+    // CHECK: function\d\(local\d\)
     _ = i();
 
-    // CHECK: j\(j\(42, local\d\), local\d\)
+    // CHECK: function\d\(function\d\(42, local\d\), local\d\)
     _ = j(j(42));
 }

Source/WebGPU/WGSL/tests/valid/name-mangling.wgsl

@@ -0,0 +1,32 @@
+// RUN: %metal myComputeEntrypoint | %check
+
+// CHECK-NOT-L: MyStruct1
+struct MyStruct1 {
+  // FIXME: mangle struct fields
+  myStructField1: i32,
+};
+
+// CHECK-NOT-L: MyStruct2
+struct MyStruct2 {
+  // FIXME: mangle struct fields
+  myStructField2: MyStruct1,
+};
+
+// FIXME: this still appears in a struct field
+@group(0) @binding(0) var<storage> myGlobal1: MyStruct2;
+
+// CHECK-NOT-L: myGlobal2
+var<private> myGlobal2: MyStruct2;
+
+// CHECK-NOT-L: myHelperFunction
+fn myHelperFunction() -> i32
+{
+  return myGlobal1.myStructField2.myStructField1 + myGlobal2.myStructField2.myStructField1;
+}
+
+// CHECK-NOT-L: myComputeEntrypoint
+@compute @workgroup_size(1)
+fn myComputeEntrypoint()
+{
+  _ = myHelperFunction();
+}

Source/WebGPU/WGSL/tests/valid/var-initialization-with-var.wgsl

@@ -10,6 +10,6 @@ fn main() {
   var b = a;
   // CHECK: local\d+ = 0
   b = 0.0;
-  // CHECK: f\(local\d+\)
+  // CHECK: function\d\(local\d+\)
   _ = f(b);
 }