From 74ef445bcafbe77eb04883b7468c619f1e1f859e Mon Sep 17 00:00:00 2001
From: Xiaopeng Li <lixiaopeng.jetspark@bytedance.com>
Date: Thu, 17 Oct 2024 17:45:29 +0800
Subject: [PATCH] ADD metal namespace for uint4 (#6417)

---
 naga/src/back/msl/writer.rs                  | 9 ++++++---
 naga/tests/out/msl/subgroup-operations-s.msl | 4 ++--
 naga/tests/out/msl/subgroup-operations.msl   | 4 ++--
 3 files changed, 10 insertions(+), 7 deletions(-)

diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs
index 7fb30b4f7..06d23d014 100644
--- a/naga/src/back/msl/writer.rs
+++ b/naga/src/back/msl/writer.rs
@@ -3323,7 +3323,10 @@ impl<W: Write> Writer<W> {
                     let name = self.namer.call("");
                     self.start_baking_expression(result, &context.expression, &name)?;
                     self.named_expressions.insert(result, name);
-                    write!(self.out, "uint4((uint64_t){NAMESPACE}::simd_ballot(")?;
+                    write!(
+                        self.out,
+                        "{NAMESPACE}::uint4((uint64_t){NAMESPACE}::simd_ballot("
+                    )?;
                     if let Some(predicate) = predicate {
                         self.put_expression(predicate, &context.expression, true)?;
                     } else {
@@ -4487,7 +4490,7 @@ template <typename A>
                 let name = self.namer.call("unpackUint32x4");
                 writeln!(
                     self.out,
-                    "uint4 {name}(uint b0, \
+                    "{NAMESPACE}::uint4 {name}(uint b0, \
                                   uint b1, \
                                   uint b2, \
                                   uint b3, \
@@ -4506,7 +4509,7 @@ template <typename A>
                 )?;
                 writeln!(
                     self.out,
-                    "{}return uint4((b3 << 24 | b2 << 16 | b1 << 8 | b0), \
+                    "{}return {NAMESPACE}::uint4((b3 << 24 | b2 << 16 | b1 << 8 | b0), \
                                     (b7 << 24 | b6 << 16 | b5 << 8 | b4), \
                                     (b11 << 24 | b10 << 16 | b9 << 8 | b8), \
                                     (b15 << 24 | b14 << 16 | b13 << 8 | b12));",
diff --git a/naga/tests/out/msl/subgroup-operations-s.msl b/naga/tests/out/msl/subgroup-operations-s.msl
index 3a6f30231..3c4df8e54 100644
--- a/naga/tests/out/msl/subgroup-operations-s.msl
+++ b/naga/tests/out/msl/subgroup-operations-s.msl
@@ -11,8 +11,8 @@ void main_1(
 ) {
     uint _e5 = subgroup_size_1;
     uint _e6 = subgroup_invocation_id_1;
-    metal::uint4 unnamed = uint4((uint64_t)metal::simd_ballot((_e6 & 1u) == 1u), 0, 0, 0);
-    metal::uint4 unnamed_1 = uint4((uint64_t)metal::simd_ballot(true), 0, 0, 0);
+    metal::uint4 unnamed = metal::uint4((uint64_t)metal::simd_ballot((_e6 & 1u) == 1u), 0, 0, 0);
+    metal::uint4 unnamed_1 = metal::uint4((uint64_t)metal::simd_ballot(true), 0, 0, 0);
     bool unnamed_2 = metal::simd_all(_e6 != 0u);
     bool unnamed_3 = metal::simd_any(_e6 == 0u);
     uint unnamed_4 = metal::simd_sum(_e6);
diff --git a/naga/tests/out/msl/subgroup-operations.msl b/naga/tests/out/msl/subgroup-operations.msl
index 980dea47f..20a550a71 100644
--- a/naga/tests/out/msl/subgroup-operations.msl
+++ b/naga/tests/out/msl/subgroup-operations.msl
@@ -19,8 +19,8 @@ kernel void main_(
 ) {
     const Structure sizes = { num_subgroups, subgroup_size };
     metal::simdgroup_barrier(metal::mem_flags::mem_threadgroup);
-    metal::uint4 unnamed = uint4((uint64_t)metal::simd_ballot((subgroup_invocation_id & 1u) == 1u), 0, 0, 0);
-    metal::uint4 unnamed_1 = uint4((uint64_t)metal::simd_ballot(true), 0, 0, 0);
+    metal::uint4 unnamed = metal::uint4((uint64_t)metal::simd_ballot((subgroup_invocation_id & 1u) == 1u), 0, 0, 0);
+    metal::uint4 unnamed_1 = metal::uint4((uint64_t)metal::simd_ballot(true), 0, 0, 0);
     bool unnamed_2 = metal::simd_all(subgroup_invocation_id != 0u);
     bool unnamed_3 = metal::simd_any(subgroup_invocation_id == 0u);
     uint unnamed_4 = metal::simd_sum(subgroup_invocation_id);