fixed incorrect packed_int3 conversion generated for Metal

Author: 39ali

naga/src/back/msl/writer.rs

@@ -1931,6 +1931,34 @@ impl<W: Write> Writer<W> {
         Ok(())
     }
 
+    fn put_unpacked_access_chain(
+        &mut self,
+        expr_handle: Handle<crate::Expression>,
+        policy: index::BoundsCheckPolicy,
+        context: &ExpressionContext,
+    ) -> Result<(), Error> {
+        let packed_kind = context.get_packed_vec_kind(expr_handle);
+        let ty = context.resolve_type(expr_handle);
+
+        // handle packed_type3 like packed_int3, we should unpack by using metal::type3()
+        if let (
+            Some(_),
+            &crate::TypeInner::Vector {
+                size: crate::VectorSize::Tri,
+                scalar,
+            },
+        ) = (packed_kind, ty)
+        {
+            put_numeric_type(&mut self.out, scalar, &[crate::VectorSize::Tri])?;
+            write!(self.out, "(")?;
+            self.put_access_chain(expr_handle, policy, context)?;
+            write!(self.out, ")")?;
+        } else {
+            self.put_access_chain(expr_handle, policy, context)?;
+        }
+        Ok(())
+    }
+
     /// Emit code for the expression `expr_handle`.
     ///
     /// The `is_scoped` argument is true if the surrounding operators have the
@@ -1993,7 +2021,7 @@ impl<W: Write> Writer<W> {
                     )?
                 {
                     write!(self.out, " ? ")?;
-                    self.put_access_chain(expr_handle, policy, context)?;
+                    self.put_unpacked_access_chain(expr_handle, policy, context)?;
                     write!(self.out, " : ")?;
 
                     if context.resolve_type(base).pointer_space().is_some() {
@@ -2035,7 +2063,7 @@ impl<W: Write> Writer<W> {
                         write!(self.out, ")")?;
                     }
                 } else {
-                    self.put_access_chain(expr_handle, policy, context)?;
+                    self.put_unpacked_access_chain(expr_handle, policy, context)?;
                 }
             }
             crate::Expression::Swizzle {

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-globals.metal

@@ -49,13 +49,13 @@ void test_msl_packed_vec3_(
     int _e16 = idx;
     alignment.v3_[_e16] = 3.0;
     FooStruct data = alignment;
-    metal::float3 l0_ = data.v3_;
-    metal::float2 l1_ = metal::float3(data.v3_).zx;
-    test_msl_packed_vec3_as_arg(data.v3_);
-    metal::float3 mvm0_ = metal::float3(data.v3_) * metal::float3x3 {};
-    metal::float3 mvm1_ = metal::float3x3 {} * metal::float3(data.v3_);
-    metal::float3 svm0_ = data.v3_ * 2.0;
-    metal::float3 svm1_ = 2.0 * data.v3_;
+    metal::float3 l0_ = metal::float3(data.v3_);
+    metal::float2 l1_ = metal::float3(metal::float3(data.v3_)).zx;
+    test_msl_packed_vec3_as_arg(metal::float3(data.v3_));
+    metal::float3 mvm0_ = metal::float3(metal::float3(data.v3_)) * metal::float3x3 {};
+    metal::float3 mvm1_ = metal::float3x3 {} * metal::float3(metal::float3(data.v3_));
+    metal::float3 svm0_ = metal::float3(data.v3_) * 2.0;
+    metal::float3 svm1_ = 2.0 * metal::float3(data.v3_);
     return;
 }
 

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 };
+}