fix(naga): zero initialize loop-local var declarations in each iteration

CHANGELOG.md

@@ -215,6 +215,7 @@ By @beholdnec in [#8505](https://github.com/gfx-rs/wgpu/pull/8505).
 - Fix typo in `naga::back::msl::Error::UnsupportedWritable*` variant names. By @ErichDonGubler in [#9376](https://github.com/gfx-rs/wgpu/pull/9376).
 - Added support for `enable wgpu_binding_array;`. By @39ali in [#9298](https://github.com/gfx-rs/wgpu/pull/9298).
 - [hlsl] more `matCx2` fixes. By @teoxoy in [#9507](https://github.com/gfx-rs/wgpu/pull/9507).
+- Fixed WGSL loop-local `var` declarations without explicit initializers so they are zero-initialized each iteration. By @ruihe774.
 
 #### Vulkan
 

naga/src/front/wgsl/lower/mod.rs

@@ -1916,6 +1916,17 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
                     let handle = ctx
                         .as_expression(block, &mut emitter)
                         .interrupt_emitter(ir::Expression::LocalVariable(var), Span::UNDEFINED)?;
+                    let initializer = if is_inside_loop {
+                        match initializer {
+                            Some(initializer) => Some(initializer),
+                            None => Some(
+                                ctx.as_expression(block, &mut emitter)
+                                    .append_expression(ir::Expression::ZeroValue(ty), stmt.span)?,
+                            ),
+                        }
+                    } else {
+                        initializer
+                    };
                     block.extend(emitter.finish(&ctx.function.expressions));
                     ctx.local_table
                         .insert(v.handle, Declared::Runtime(Typed::Reference(handle)));

naga/tests/in/wgsl/9489-loop-local-var-init.toml

@@ -0,0 +1,11 @@
+targets = "SPIRV | METAL | GLSL | HLSL | WGSL | IR"
+
+[msl.per_entry_point_map.main]
+resources = [
+    { bind_target = { buffer = 0, mutable = false }, resource_binding = { group = 0, binding = 0 } },
+    { bind_target = { buffer = 1, mutable = true }, resource_binding = { group = 0, binding = 1 } },
+]
+sizes_buffer = 24
+
+[glsl]
+version.Desktop = 430

naga/tests/in/wgsl/9489-loop-local-var-init.wgsl

@@ -0,0 +1,30 @@
+// #9489: A `var` declaration inside a loop body without an explicit
+// initializer must be re-zero-initialized on every iteration, just like
+// one with an explicit initializer. Naga hoists all local variables to
+// function scope and zero-initializes them once at function entry, so the
+// per-iteration reset has to be lowered to an explicit store in the loop
+// body (the same mechanism already used for explicit initializers).
+//
+// Without the fix, `acc_noinit` accumulates across iterations of the outer
+// loop (a running prefix sum) while `acc_init` does not, even though the
+// two are semantically equivalent per the WGSL spec.
+
+@group(0) @binding(0) var<storage, read> input: array<f32, 64>;
+@group(0) @binding(1) var<storage, read_write> output: array<f32, 8>;
+
+@compute @workgroup_size(1)
+fn main() {
+    for (var t = 0u; t < 4u; t++) {
+        var acc_noinit: vec4<f32>;                // no explicit initializer
+        var acc_init: vec4<f32> = vec4<f32>();    // explicit initializer
+
+        for (var d = 0u; d < 16u; d++) {
+            let v = vec4<f32>(input[t * 16u + d]);
+            acc_noinit += v;
+            acc_init += v;
+        }
+
+        output[t * 2u] = acc_noinit.x;
+        output[t * 2u + 1u] = acc_init.x;
+    }
+}

naga/tests/out/glsl/wgsl-9489-loop-local-var-init.main.Compute.glsl

@@ -0,0 +1,65 @@
+#version 430 core
+#extension GL_ARB_compute_shader : require
+#extension GL_ARB_shader_storage_buffer_object : require
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+layout(std430) readonly buffer type_1_block_0Compute { float _group_0_binding_0_cs[64]; };
+
+layout(std430) buffer type_2_block_1Compute { float _group_0_binding_1_cs[8]; };
+
+
+void main() {
+    uint t = 0u;
+    vec4 acc_noinit = vec4(0.0);
+    vec4 acc_init = vec4(0.0);
+    uint d = 0u;
+    bool loop_init = true;
+    while(true) {
+        if (!loop_init) {
+            uint _e47 = t;
+            t = (_e47 + 1u);
+        }
+        loop_init = false;
+        uint _e2 = t;
+        if ((_e2 < 4u)) {
+        } else {
+            break;
+        }
+        {
+            acc_noinit = vec4(0.0);
+            acc_init = vec4(0.0);
+            d = 0u;
+            bool loop_init_1 = true;
+            while(true) {
+                if (!loop_init_1) {
+                    uint _e28 = d;
+                    d = (_e28 + 1u);
+                }
+                loop_init_1 = false;
+                uint _e11 = d;
+                if ((_e11 < 16u)) {
+                } else {
+                    break;
+                }
+                {
+                    uint _e15 = t;
+                    uint _e18 = d;
+                    float _e21 = _group_0_binding_0_cs[((_e15 * 16u) + _e18)];
+                    vec4 v = vec4(_e21);
+                    vec4 _e23 = acc_noinit;
+                    acc_noinit = (_e23 + v);
+                    vec4 _e25 = acc_init;
+                    acc_init = (_e25 + v);
+                }
+            }
+            uint _e31 = t;
+            float _e36 = acc_noinit.x;
+            _group_0_binding_1_cs[(_e31 * 2u)] = _e36;
+            uint _e38 = t;
+            float _e45 = acc_init.x;
+            _group_0_binding_1_cs[((_e38 * 2u) + 1u)] = _e45;
+        }
+    }
+    return;
+}
+

naga/tests/out/hlsl/wgsl-9489-loop-local-var-init.hlsl

@@ -0,0 +1,70 @@
+ByteAddressBuffer input : register(t0);
+RWByteAddressBuffer output : register(u1);
+
+float4 ZeroValuefloat4() {
+    return (float4)0;
+}
+
+[numthreads(1, 1, 1)]
+void main()
+{
+    uint t = 0u;
+    float4 acc_noinit = (float4)0;
+    float4 acc_init = (float4)0;
+    uint d = (uint)0;
+
+    uint2 loop_bound = uint2(4294967295u, 4294967295u);
+    bool loop_init = true;
+    while(true) {
+        if (all(loop_bound == uint2(0u, 0u))) { break; }
+        loop_bound -= uint2(loop_bound.y == 0u, 1u);
+        if (!loop_init) {
+            uint _e47 = t;
+            t = (_e47 + 1u);
+        }
+        loop_init = false;
+        uint _e2 = t;
+        if ((_e2 < 4u)) {
+        } else {
+            break;
+        }
+        {
+            acc_noinit = ZeroValuefloat4();
+            acc_init = ZeroValuefloat4();
+            d = 0u;
+            uint2 loop_bound_1 = uint2(4294967295u, 4294967295u);
+            bool loop_init_1 = true;
+            while(true) {
+                if (all(loop_bound_1 == uint2(0u, 0u))) { break; }
+                loop_bound_1 -= uint2(loop_bound_1.y == 0u, 1u);
+                if (!loop_init_1) {
+                    uint _e28 = d;
+                    d = (_e28 + 1u);
+                }
+                loop_init_1 = false;
+                uint _e11 = d;
+                if ((_e11 < 16u)) {
+                } else {
+                    break;
+                }
+                {
+                    uint _e15 = t;
+                    uint _e18 = d;
+                    float _e21 = asfloat(input.Load(((_e15 * 16u) + _e18)*4));
+                    float4 v = (_e21).xxxx;
+                    float4 _e23 = acc_noinit;
+                    acc_noinit = (_e23 + v);
+                    float4 _e25 = acc_init;
+                    acc_init = (_e25 + v);
+                }
+            }
+            uint _e31 = t;
+            float _e36 = acc_noinit.x;
+            output.Store((_e31 * 2u)*4, asuint(_e36));
+            uint _e38 = t;
+            float _e45 = acc_init.x;
+            output.Store(((_e38 * 2u) + 1u)*4, asuint(_e45));
+        }
+    }
+    return;
+}

naga/tests/out/hlsl/wgsl-9489-loop-local-var-init.ron

@@ -0,0 +1,16 @@
+(
+    vertex:[
+    ],
+    fragment:[
+    ],
+    compute:[
+        (
+            entry_point:"main",
+            target_profile:"cs_5_1",
+        ),
+    ],
+    task:[
+    ],
+    mesh:[
+    ],
+)