From 59a2441a432e96bf174fac814cede612f9fb0c41 Mon Sep 17 00:00:00 2001 From: Jim Blandy Date: Tue, 9 May 2023 16:20:15 -0700 Subject: [PATCH] [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 | 24 +++++------------------- tests/out/msl/access.msl | 10 +++++----- tests/out/msl/policy-mix.msl | 2 +- tests/out/msl/workgroup-var-init.msl | 2 +- 4 files changed, 12 insertions(+), 26 deletions(-) diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index 88600a8f11..26db61eed7 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -2954,31 +2954,17 @@ impl Writer { 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}" diff --git a/tests/out/msl/access.msl b/tests/out/msl/access.msl index a4967117ff..4e51de80aa 100644 --- a/tests/out/msl/access.msl +++ b/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(b), 3, 4, 5}.inner[_i]; + c2_ = type_22 {a_1, static_cast(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; diff --git a/tests/out/msl/policy-mix.msl b/tests/out/msl/policy-mix.msl index 7eb0c61ede..7f270e3045 100644 --- a/tests/out/msl/policy-mix.msl +++ b/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()); diff --git a/tests/out/msl/workgroup-var-init.msl b/tests/out/msl/workgroup-var-init.msl index 0d61e96f7c..ac300d4337 100644 --- a/tests/out/msl/workgroup-var-init.msl +++ b/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; }