[msl-out] zero init variables in function address space

src/back/msl/writer.rs

@@ -3269,15 +3269,20 @@ impl<W: Write> Writer<W> {
                 };
                 let local_name = &self.names[&NameKey::FunctionLocal(fun_handle, local_handle)];
                 write!(self.out, "{}{} {}", back::INDENT, ty_name, local_name)?;
-                if let Some(value) = local.init {
-                    let coco = ConstantContext {
-                        handle: value,
-                        arena: &module.constants,
-                        names: &self.names,
-                        first_time: false,
-                    };
-                    write!(self.out, " = {}", coco)?;
-                }
+                match local.init {
+                    Some(value) => {
+                        let coco = ConstantContext {
+                            handle: value,
+                            arena: &module.constants,
+                            names: &self.names,
+                            first_time: false,
+                        };
+                        write!(self.out, " = {}", coco)?;
+                    }
+                    None => {
+                        write!(self.out, " = {{}}")?;
+                    }
+                };
                 writeln!(self.out, ";")?;
             }
 
@@ -3801,15 +3806,20 @@ impl<W: Write> Writer<W> {
                     first_time: false,
                 };
                 write!(self.out, "{}{} {}", back::INDENT, ty_name, name)?;
-                if let Some(value) = local.init {
-                    let coco = ConstantContext {
-                        handle: value,
-                        arena: &module.constants,
-                        names: &self.names,
-                        first_time: false,
-                    };
-                    write!(self.out, " = {}", coco)?;
-                }
+                match local.init {
+                    Some(value) => {
+                        let coco = ConstantContext {
+                            handle: value,
+                            arena: &module.constants,
+                            names: &self.names,
+                            first_time: false,
+                        };
+                        write!(self.out, " = {}", coco)?;
+                    }
+                    None => {
+                        write!(self.out, " = {{}}")?;
+                    }
+                };
                 writeln!(self.out, ";")?;
             }
 

tests/out/msl/access.msl

@@ -45,7 +45,7 @@ void test_matrix_within_struct_accesses(
     constant Baz& baz
 ) {
     int idx = 9;
-    Baz t;
+    Baz t = {};
     int _e4 = idx;
     idx = _e4 - 1;
     metal::float3x2 unnamed = baz.m;
@@ -103,7 +103,7 @@ vertex foo_vertOutput foo_vert(
 , constant _mslBufferSizes& _buffer_sizes [[buffer(24)]]
 ) {
     float foo = 0.0;
-    type_16 c;
+    type_16 c = {};
     float baz_1 = foo;
     foo = 1.0;
     test_matrix_within_struct_accesses(baz);
@@ -139,7 +139,7 @@ kernel void atomics(
   device Bar& bar [[buffer(0)]]
 , constant _mslBufferSizes& _buffer_sizes [[buffer(24)]]
 ) {
-    int tmp;
+    int tmp = {};
     int value_1 = metal::atomic_load_explicit(&bar.atom, metal::memory_order_relaxed);
     int _e7 = metal::atomic_fetch_add_explicit(&bar.atom, 5, metal::memory_order_relaxed);
     tmp = _e7;

tests/out/msl/binding-arrays.msl

@@ -37,9 +37,9 @@ fragment main_Output main_(
 ) {
     const FragmentIn fragment_in = { varyings.index };
     int i1_ = 0;
-    metal::int2 i2_;
+    metal::int2 i2_ = {};
     float v1_ = 0.0;
-    metal::float4 v4_;
+    metal::float4 v4_ = {};
     uint uniform_index = uni.index;
     uint non_uniform_index = fragment_in.index;
     i2_ = metal::int2(0);

tests/out/msl/bits.msl

@@ -8,15 +8,15 @@ using metal::uint;
 kernel void main_(
 ) {
     int i = 0;
-    metal::int2 i2_;
-    metal::int3 i3_;
-    metal::int4 i4_;
+    metal::int2 i2_ = {};
+    metal::int3 i3_ = {};
+    metal::int4 i4_ = {};
     uint u = 0u;
-    metal::uint2 u2_;
-    metal::uint3 u3_;
-    metal::uint4 u4_;
-    metal::float2 f2_;
-    metal::float4 f4_;
+    metal::uint2 u2_ = {};
+    metal::uint3 u3_ = {};
+    metal::uint4 u4_ = {};
+    metal::float2 f2_ = {};
+    metal::float4 f4_ = {};
     i2_ = metal::int2(0);
     i3_ = metal::int3(0);
     i4_ = metal::int4(0);

tests/out/msl/boids.msl

@@ -37,15 +37,15 @@ kernel void main_(
 , device Particles& particlesDst [[buffer(2)]]
 , constant _mslBufferSizes& _buffer_sizes [[buffer(3)]]
 ) {
-    metal::float2 vPos;
-    metal::float2 vVel;
-    metal::float2 cMass;
-    metal::float2 cVel;
-    metal::float2 colVel;
+    metal::float2 vPos = {};
+    metal::float2 vVel = {};
+    metal::float2 cMass = {};
+    metal::float2 cVel = {};
+    metal::float2 colVel = {};
     int cMassCount = 0;
     int cVelCount = 0;
-    metal::float2 pos;
-    metal::float2 vel;
+    metal::float2 pos = {};
+    metal::float2 vel = {};
     uint i = 0u;
     uint index = global_invocation_id.x;
     if (index >= NUM_PARTICLES) {

tests/out/msl/collatz.msl

@@ -16,7 +16,7 @@ struct PrimeIndices {
 uint collatz_iterations(
     uint n_base
 ) {
-    uint n;
+    uint n = {};
     uint i = 0u;
     n = n_base;
     while(true) {

tests/out/msl/control-flow.msl

@@ -49,7 +49,7 @@ struct main_Input {
 kernel void main_(
   metal::uint3 global_id [[thread_position_in_grid]]
 ) {
-    int pos;
+    int pos = {};
     metal::threadgroup_barrier(metal::mem_flags::mem_device);
     metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
     switch(1) {

tests/out/msl/interpolate.msl

@@ -27,7 +27,7 @@ struct vert_mainOutput {
 };
 vertex vert_mainOutput vert_main(
 ) {
-    FragmentInput out;
+    FragmentInput out = {};
     out.position = metal::float4(2.0, 4.0, 5.0, 6.0);
     out._flat = 8u;
     out._linear = 27.0;

tests/out/msl/operators.msl

@@ -58,7 +58,7 @@ metal::float3 bool_cast(
 
 float constructors(
 ) {
-    Foo foo;
+    Foo foo = {};
     foo = Foo {metal::float4(1.0), 1};
     metal::float2x2 mat2comp = metal::float2x2(metal::float2(1.0, 0.0), metal::float2(0.0, 1.0));
     metal::float4x4 mat4comp = metal::float4x4(metal::float4(1.0, 0.0, 0.0, 0.0), metal::float4(0.0, 1.0, 0.0, 0.0), metal::float4(0.0, 0.0, 1.0, 0.0), metal::float4(0.0, 0.0, 0.0, 1.0));

tests/out/msl/policy-mix.msl

@@ -42,7 +42,7 @@ metal::float4 mock_function(
     threadgroup type_5 const& in_workgroup,
     thread type_6 const& in_private
 ) {
-    type_9 in_function;
+    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];
     metal::float4 _e22 = in_storage.a.inner[i];
     metal::float4 _e25 = in_uniform.a.inner[i];

tests/out/msl/shadow.msl

@@ -65,7 +65,7 @@ vertex vs_mainOutput vs_main(
 ) {
     const auto position = varyings.position;
     const auto normal = varyings.normal;
-    VertexOutput out;
+    VertexOutput out = {};
     metal::float4x4 w = u_entity.world;
     metal::float4x4 _e7 = u_entity.world;
     metal::float4 world_pos = _e7 * static_cast<metal::float4>(position);

tests/out/msl/skybox.msl

@@ -23,8 +23,8 @@ vertex vs_mainOutput vs_main(
   uint vertex_index [[vertex_id]]
 , constant Data& r_data [[buffer(0)]]
 ) {
-    int tmp1_;
-    int tmp2_;
+    int tmp1_ = {};
+    int tmp2_ = {};
     tmp1_ = static_cast<int>(vertex_index) / 2;
     tmp2_ = static_cast<int>(vertex_index) & 1;
     int _e10 = tmp1_;