fixed incorrect packed_int3 conversion generated for Metal

Author: 39ali

naga/src/back/msl/writer.rs

@@ -2274,6 +2274,7 @@ impl<W: Write> Writer<W> {
                     // to signed.
                     self.put_bitcasted_expression(
                         context.resolve_type(expr_handle),
+                        expr_handle,
                         context,
                         &|writer, context, is_scoped| {
                             writer.put_binop(
@@ -2285,6 +2286,7 @@ impl<W: Write> Writer<W> {
                                 &|writer, expr, context, _is_scoped| {
                                     writer.put_bitcasted_expression(
                                         &to_unsigned(context.resolve_type(expr))?,
+                                        expr,
                                         context,
                                         &|writer, context, is_scoped| {
                                             writer.put_expression(expr, context, is_scoped)
@@ -3026,6 +3028,7 @@ impl<W: Write> Writer<W> {
     fn put_bitcasted_expression<F>(
         &mut self,
         cast_to: &crate::TypeInner,
+        inner_expr: Handle<crate::Expression>,
         context: &ExpressionContext,
         put_expression: &F,
     ) -> BackendResult
@@ -3041,9 +3044,18 @@ impl<W: Write> Writer<W> {
             _ => return Err(Error::UnsupportedBitCast(cast_to.clone())),
         };
         write!(self.out, ">(")?;
-        put_expression(self, context, true)?;
-        write!(self.out, ")")?;
 
+        // if it's packed, we must unpack it (e.g., float3(val)) before the bitcast.
+        if let Some(scalar) = context.get_packed_vec_kind(inner_expr) {
+            put_numeric_type(&mut self.out, scalar, &[crate::VectorSize::Tri])?;
+            write!(self.out, "(")?;
+            put_expression(self, context, true)?;
+            write!(self.out, ")")?;
+        } else {
+            put_expression(self, context, true)?;
+        }
+
+        write!(self.out, ")")?;
         Ok(())
     }
 

naga/tests/in/wgsl/packed-vec3-bitcast.toml

@@ -0,0 +1 @@
+targets = "METAL"

naga/tests/in/wgsl/packed-vec3-bitcast.wgsl

@@ -0,0 +1,19 @@
+struct VertexInput {
+    @location(0) chunk : vec3<i32>,
+    @location(1) texture_index : u32,
+}
+
+struct VertexOutput {
+    @builtin(position) clip_position : vec4<f32>,
+};
+
+@vertex
+fn vs_main(
+in : VertexInput,
+) -> VertexOutput {
+    var out : VertexOutput;
+
+    let position = vec3<f32> (in.chunk - vec3<i32>(5));
+
+    return out;
+}

naga/tests/out/msl/wgsl-packed-vec3-bitcast.metal

@@ -0,0 +1,31 @@
+// language: metal1.0
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using metal::uint;
+
+struct VertexInput {
+    metal::packed_int3 chunk;
+    uint texture_index;
+};
+struct VertexOutput {
+    metal::float4 clip_position;
+};
+
+struct vs_mainInput {
+    metal::int3 chunk [[attribute(0)]];
+    uint texture_index [[attribute(1)]];
+};
+struct vs_mainOutput {
+    metal::float4 clip_position [[position]];
+};
+vertex vs_mainOutput vs_main(
+  vs_mainInput varyings [[stage_in]]
+) {
+    const VertexInput in = { varyings.chunk, varyings.texture_index };
+    VertexOutput out = {};
+    metal::float3 position = static_cast<metal::float3>(as_type<metal::int3>(as_type<metal::uint3>(metal::int3(in.chunk)) - as_type<metal::uint3>(metal::int3(5))));
+    VertexOutput _e7 = out;
+    const auto _tmp = _e7;
+    return vs_mainOutput { _tmp.clip_position };
+}