Test

gfx-rs · Sep 30, 2024 · 631640d · 631640d
1 parent 2662c06
commit 631640d
Show file tree

Hide file tree

Showing 2 changed files with 4 additions and 4 deletions.
diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs
@@ -1205,9 +1205,9 @@ impl<W: Write> Writer<W> {
         write!(self.out, ".atomic_{}(", fun.to_msl_64_bit()?)?;
         // coordinates in IR are int, but Metal expects uint
         self.put_cast_to_uint_scalar_or_vector(address.coordinate, &context.expression)?;
-        write!(self.out, ", ")?;
+        write!(self.out, ",  metal::ulong4(")?;
         self.put_expression(value, &context.expression, true)?;
-        writeln!(self.out, ");")?;
+        writeln!(self.out, ", 0uL, 0uL, 0uL));")?;
 
         Ok(())
     }

diff --git a/naga/tests/out/msl/atomicTexture-int64.msl b/naga/tests/out/msl/atomicTexture-int64.msl
@@ -11,8 +11,8 @@ kernel void cs_main(
   metal::uint3 id [[thread_position_in_threadgroup]]
 , metal::texture2d<ulong, metal::access::write> image [[user(fake0)]]
 ) {
-    image.atomic_max(metal::uint2(metal::int2(0, 0)), 1uL);
+    image.atomic_max(metal::uint2(metal::int2(0, 0)),  metal::ulong4(1uL, 0uL, 0uL, 0uL));
     metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
-    image.atomic_min(metal::uint2(metal::int2(0, 0)), 1uL);
+    image.atomic_min(metal::uint2(metal::int2(0, 0)),  metal::ulong4(1uL, 0uL, 0uL, 0uL));
     return;
 }