[msl-out] Revert af4d989 since #764 removes the need for it (#2331)

Assigning arrays by value works fine since all arrays are now wrapped by a struct.

Co-authored-by: teoxoy <28601907+teoxoy@users.noreply.github.com>

src/back/msl/writer.rs

@@ -2954,31 +2954,17 @@ impl<W: Write> Writer<W> {
         context: &StatementContext,
     ) -> BackendResult {
         let pointer_inner = context.expression.resolve_type(pointer);
-        let (array_size, is_atomic) = match *pointer_inner {
+        let is_atomic = match *pointer_inner {
             crate::TypeInner::Pointer { base, .. } => {
                 match context.expression.module.types[base].inner {
-                    crate::TypeInner::Array {
-                        size: crate::ArraySize::Constant(ch),
-                        ..
-                    } => (Some(ch), false),
-                    crate::TypeInner::Atomic { .. } => (None, true),
-                    _ => (None, false),
+                    crate::TypeInner::Atomic { .. } => true,
+                    _ => false,
                 }
             }
-            _ => (None, false),
+            _ => false,
         };
 
-        // we can't assign fixed-size arrays
-        if let Some(const_handle) = array_size {
-            let size = context.expression.module.constants[const_handle]
-                .to_array_length()
-                .unwrap();
-            write!(self.out, "{level}for(int _i=0; _i<{size}; ++_i) ")?;
-            self.put_access_chain(pointer, policy, &context.expression)?;
-            write!(self.out, ".{WRAPPED_ARRAY_FIELD}[_i] = ")?;
-            self.put_expression(value, &context.expression, true)?;
-            writeln!(self.out, ".{WRAPPED_ARRAY_FIELD}[_i];")?;
-        } else if is_atomic {
+        if is_atomic {
             write!(
                 self.out,
                 "{level}{NAMESPACE}::atomic_store_explicit({ATOMIC_REFERENCE}"

tests/out/msl/access.msl

@@ -128,7 +128,7 @@ void test_matrix_within_array_within_struct_accesses(
     t_1 = MatCx2InArray {const_type_15_};
     int _e66 = idx_1;
     idx_1 = _e66 + 1;
-    for(int _i=0; _i<2; ++_i) t_1.am.inner[_i] = const_type_15_.inner[_i];
+    t_1.am = const_type_15_;
     t_1.am.inner[0] = metal::float4x2(metal::float2(8.0), metal::float2(7.0), metal::float2(6.0), metal::float2(5.0));
     t_1.am.inner[0][0] = metal::float2(9.0);
     int _e93 = idx_1;
@@ -167,7 +167,7 @@ void assign_through_ptr_fn(
 void assign_array_through_ptr_fn(
     thread type_26& foo_2
 ) {
-    for(int _i=0; _i<2; ++_i) foo_2.inner[_i] = type_26 {metal::float4(1.0), metal::float4(2.0)}.inner[_i];
+    foo_2 = type_26 {metal::float4(1.0), metal::float4(2.0)};
     return;
 }
 
@@ -197,7 +197,7 @@ vertex foo_vertOutput foo_vert(
     int a_1 = bar.data[(1 + (_buffer_sizes.size1 - 160 - 8) / 8) - 2u].value;
     metal::int2 c = qux;
     float _e34 = read_from_private(foo);
-    for(int _i=0; _i<5; ++_i) c2_.inner[_i] = type_22 {a_1, static_cast<int>(b), 3, 4, 5}.inner[_i];
+    c2_ = type_22 {a_1, static_cast<int>(b), 3, 4, 5};
     c2_.inner[vi + 1u] = 42;
     int value = c2_.inner[vi];
     float _e48 = test_arr_as_arg(const_type_19_);
@@ -215,7 +215,7 @@ fragment foo_fragOutput foo_frag(
 ) {
     bar._matrix[1].z = 1.0;
     bar._matrix = metal::float4x3(metal::float3(0.0), metal::float3(1.0), metal::float3(2.0), metal::float3(3.0));
-    for(int _i=0; _i<2; ++_i) bar.arr.inner[_i] = type_9 {metal::uint2(0u), metal::uint2(1u)}.inner[_i];
+    bar.arr = type_9 {metal::uint2(0u), metal::uint2(1u)};
     bar.data[1].value = 1;
     qux = const_type_12_;
     return foo_fragOutput { metal::float4(0.0) };
@@ -231,7 +231,7 @@ kernel void assign_through_ptr(
     }
     metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
     type_26 arr = {};
-    for(int _i=0; _i<2; ++_i) arr.inner[_i] = type_26 {metal::float4(6.0), metal::float4(7.0)}.inner[_i];
+    arr = type_26 {metal::float4(6.0), metal::float4(7.0)};
     assign_through_ptr_fn(val);
     assign_array_through_ptr_fn(arr);
     return;

tests/out/msl/policy-mix.msl

@@ -43,7 +43,7 @@ metal::float4 mock_function(
     thread type_6& in_private
 ) {
     type_9 in_function = {};
-    for(int _i=0; _i<2; ++_i) in_function.inner[_i] = type_9 {metal::float4(0.7070000171661377, 0.0, 0.0, 1.0), metal::float4(0.0, 0.7070000171661377, 0.0, 1.0)}.inner[_i];
+    in_function = type_9 {metal::float4(0.7070000171661377, 0.0, 0.0, 1.0), metal::float4(0.0, 0.7070000171661377, 0.0, 1.0)};
     metal::float4 _e18 = in_storage.a.inner[i];
     metal::float4 _e22 = in_uniform.a.inner[i];
     metal::float4 _e25 = (uint(l) < image_2d_array.get_num_mip_levels() && uint(i) < image_2d_array.get_array_size() && metal::all(metal::uint2(c) < metal::uint2(image_2d_array.get_width(l), image_2d_array.get_height(l))) ? image_2d_array.read(metal::uint2(c), i, l): DefaultConstructible());

tests/out/msl/workgroup-var-init.msl

@@ -35,6 +35,6 @@ kernel void main_(
     }
     metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
     type_1 _e3 = w_mem.arr;
-    for(int _i=0; _i<512; ++_i) output.inner[_i] = _e3.inner[_i];
+    output = _e3;
     return;
 }