diff --git a/CHANGELOG.md b/CHANGELOG.md index 24f9c31a726..a235eb0d8e5 100644 --- a/CHANGELOG.md +++ b/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 diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs index 2572b07c122..0d65416009f 100644 --- a/naga/src/front/wgsl/lower/mod.rs +++ b/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))); diff --git a/naga/tests/in/wgsl/9489-loop-local-var-init.toml b/naga/tests/in/wgsl/9489-loop-local-var-init.toml new file mode 100644 index 00000000000..fcf5add6000 --- /dev/null +++ b/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 diff --git a/naga/tests/in/wgsl/9489-loop-local-var-init.wgsl b/naga/tests/in/wgsl/9489-loop-local-var-init.wgsl new file mode 100644 index 00000000000..ad6f6ce893f --- /dev/null +++ b/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 input: array; +@group(0) @binding(1) var output: array; + +@compute @workgroup_size(1) +fn main() { + for (var t = 0u; t < 4u; t++) { + var acc_noinit: vec4; // no explicit initializer + var acc_init: vec4 = vec4(); // explicit initializer + + for (var d = 0u; d < 16u; d++) { + let v = vec4(input[t * 16u + d]); + acc_noinit += v; + acc_init += v; + } + + output[t * 2u] = acc_noinit.x; + output[t * 2u + 1u] = acc_init.x; + } +} diff --git a/naga/tests/out/glsl/wgsl-9489-loop-local-var-init.main.Compute.glsl b/naga/tests/out/glsl/wgsl-9489-loop-local-var-init.main.Compute.glsl new file mode 100644 index 00000000000..1e1c96a87c7 --- /dev/null +++ b/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; +} + diff --git a/naga/tests/out/hlsl/wgsl-9489-loop-local-var-init.hlsl b/naga/tests/out/hlsl/wgsl-9489-loop-local-var-init.hlsl new file mode 100644 index 00000000000..aa88c6e0fc5 --- /dev/null +++ b/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; +} diff --git a/naga/tests/out/hlsl/wgsl-9489-loop-local-var-init.ron b/naga/tests/out/hlsl/wgsl-9489-loop-local-var-init.ron new file mode 100644 index 00000000000..f1edcbbbdd1 --- /dev/null +++ b/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:[ + ], +) diff --git a/naga/tests/out/ir/wgsl-9489-loop-local-var-init.compact.ron b/naga/tests/out/ir/wgsl-9489-loop-local-var-init.compact.ron new file mode 100644 index 00000000000..fc649ff2c71 --- /dev/null +++ b/naga/tests/out/ir/wgsl-9489-loop-local-var-init.compact.ron @@ -0,0 +1,402 @@ +( + types: [ + ( + name: None, + inner: Scalar(( + kind: Float, + width: 4, + )), + ), + ( + name: None, + inner: Array( + base: 0, + size: Constant(64), + stride: 4, + ), + ), + ( + name: None, + inner: Array( + base: 0, + size: Constant(8), + stride: 4, + ), + ), + ( + name: None, + inner: Scalar(( + kind: Uint, + width: 4, + )), + ), + ( + name: None, + inner: Vector( + size: Quad, + scalar: ( + kind: Float, + width: 4, + ), + ), + ), + ], + special_types: ( + ray_desc: None, + ray_intersection: None, + ray_vertex_return: None, + external_texture_params: None, + external_texture_transfer_function: None, + predeclared_types: {}, + ), + constants: [], + overrides: [], + global_variables: [ + ( + name: Some("input"), + space: Storage( + access: ("LOAD"), + ), + binding: Some(( + group: 0, + binding: 0, + )), + ty: 1, + init: None, + memory_decorations: (""), + ), + ( + name: Some("output"), + space: Storage( + access: ("LOAD | STORE"), + ), + binding: Some(( + group: 0, + binding: 1, + )), + ty: 2, + init: None, + memory_decorations: (""), + ), + ], + global_expressions: [], + functions: [], + entry_points: [ + ( + name: "main", + stage: Compute, + early_depth_test: None, + workgroup_size: (1, 1, 1), + workgroup_size_overrides: None, + function: ( + name: Some("main"), + arguments: [], + result: None, + local_variables: [ + ( + name: Some("t"), + ty: 3, + init: Some(0), + ), + ( + name: Some("acc_noinit"), + ty: 4, + init: None, + ), + ( + name: Some("acc_init"), + ty: 4, + init: None, + ), + ( + name: Some("d"), + ty: 3, + init: None, + ), + ], + expressions: [ + Literal(U32(0)), + LocalVariable(0), + Load( + pointer: 1, + ), + Literal(U32(4)), + Binary( + op: Less, + left: 2, + right: 3, + ), + LocalVariable(1), + ZeroValue(4), + ZeroValue(4), + LocalVariable(2), + Literal(U32(0)), + LocalVariable(3), + Load( + pointer: 10, + ), + Literal(U32(16)), + Binary( + op: Less, + left: 11, + right: 12, + ), + GlobalVariable(0), + Load( + pointer: 1, + ), + Literal(U32(16)), + Binary( + op: Multiply, + left: 15, + right: 16, + ), + Load( + pointer: 10, + ), + Binary( + op: Add, + left: 17, + right: 18, + ), + Access( + base: 14, + index: 19, + ), + Load( + pointer: 20, + ), + Splat( + size: Quad, + value: 21, + ), + Load( + pointer: 5, + ), + Binary( + op: Add, + left: 23, + right: 22, + ), + Load( + pointer: 8, + ), + Binary( + op: Add, + left: 25, + right: 22, + ), + Literal(U32(1)), + Load( + pointer: 10, + ), + Binary( + op: Add, + left: 28, + right: 27, + ), + GlobalVariable(1), + Load( + pointer: 1, + ), + Literal(U32(2)), + Binary( + op: Multiply, + left: 31, + right: 32, + ), + Access( + base: 30, + index: 33, + ), + AccessIndex( + base: 5, + index: 0, + ), + Load( + pointer: 35, + ), + GlobalVariable(1), + Load( + pointer: 1, + ), + Literal(U32(2)), + Binary( + op: Multiply, + left: 38, + right: 39, + ), + Literal(U32(1)), + Binary( + op: Add, + left: 40, + right: 41, + ), + Access( + base: 37, + index: 42, + ), + AccessIndex( + base: 8, + index: 0, + ), + Load( + pointer: 44, + ), + Literal(U32(1)), + Load( + pointer: 1, + ), + Binary( + op: Add, + left: 47, + right: 46, + ), + ], + named_expressions: { + 22: "v", + }, + body: [ + Loop( + body: [ + Emit(( + start: 2, + end: 3, + )), + Emit(( + start: 4, + end: 5, + )), + If( + condition: 4, + accept: [], + reject: [ + Break, + ], + ), + Block([ + Store( + pointer: 5, + value: 6, + ), + Store( + pointer: 8, + value: 7, + ), + Store( + pointer: 10, + value: 9, + ), + Loop( + body: [ + Emit(( + start: 11, + end: 12, + )), + Emit(( + start: 13, + end: 14, + )), + If( + condition: 13, + accept: [], + reject: [ + Break, + ], + ), + Block([ + Emit(( + start: 15, + end: 16, + )), + Emit(( + start: 17, + end: 23, + )), + Emit(( + start: 23, + end: 25, + )), + Store( + pointer: 5, + value: 24, + ), + Emit(( + start: 25, + end: 27, + )), + Store( + pointer: 8, + value: 26, + ), + ]), + ], + continuing: [ + Emit(( + start: 28, + end: 30, + )), + Store( + pointer: 10, + value: 29, + ), + ], + break_if: None, + ), + Emit(( + start: 31, + end: 32, + )), + Emit(( + start: 33, + end: 37, + )), + Store( + pointer: 34, + value: 36, + ), + Emit(( + start: 38, + end: 39, + )), + Emit(( + start: 40, + end: 41, + )), + Emit(( + start: 42, + end: 46, + )), + Store( + pointer: 43, + value: 45, + ), + ]), + ], + continuing: [ + Emit(( + start: 47, + end: 49, + )), + Store( + pointer: 1, + value: 48, + ), + ], + break_if: None, + ), + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + mesh_info: None, + task_payload: None, + incoming_ray_payload: None, + ), + ], + diagnostic_filters: [], + diagnostic_filter_leaf: None, + doc_comments: None, +) \ No newline at end of file diff --git a/naga/tests/out/ir/wgsl-9489-loop-local-var-init.ron b/naga/tests/out/ir/wgsl-9489-loop-local-var-init.ron new file mode 100644 index 00000000000..fc649ff2c71 --- /dev/null +++ b/naga/tests/out/ir/wgsl-9489-loop-local-var-init.ron @@ -0,0 +1,402 @@ +( + types: [ + ( + name: None, + inner: Scalar(( + kind: Float, + width: 4, + )), + ), + ( + name: None, + inner: Array( + base: 0, + size: Constant(64), + stride: 4, + ), + ), + ( + name: None, + inner: Array( + base: 0, + size: Constant(8), + stride: 4, + ), + ), + ( + name: None, + inner: Scalar(( + kind: Uint, + width: 4, + )), + ), + ( + name: None, + inner: Vector( + size: Quad, + scalar: ( + kind: Float, + width: 4, + ), + ), + ), + ], + special_types: ( + ray_desc: None, + ray_intersection: None, + ray_vertex_return: None, + external_texture_params: None, + external_texture_transfer_function: None, + predeclared_types: {}, + ), + constants: [], + overrides: [], + global_variables: [ + ( + name: Some("input"), + space: Storage( + access: ("LOAD"), + ), + binding: Some(( + group: 0, + binding: 0, + )), + ty: 1, + init: None, + memory_decorations: (""), + ), + ( + name: Some("output"), + space: Storage( + access: ("LOAD | STORE"), + ), + binding: Some(( + group: 0, + binding: 1, + )), + ty: 2, + init: None, + memory_decorations: (""), + ), + ], + global_expressions: [], + functions: [], + entry_points: [ + ( + name: "main", + stage: Compute, + early_depth_test: None, + workgroup_size: (1, 1, 1), + workgroup_size_overrides: None, + function: ( + name: Some("main"), + arguments: [], + result: None, + local_variables: [ + ( + name: Some("t"), + ty: 3, + init: Some(0), + ), + ( + name: Some("acc_noinit"), + ty: 4, + init: None, + ), + ( + name: Some("acc_init"), + ty: 4, + init: None, + ), + ( + name: Some("d"), + ty: 3, + init: None, + ), + ], + expressions: [ + Literal(U32(0)), + LocalVariable(0), + Load( + pointer: 1, + ), + Literal(U32(4)), + Binary( + op: Less, + left: 2, + right: 3, + ), + LocalVariable(1), + ZeroValue(4), + ZeroValue(4), + LocalVariable(2), + Literal(U32(0)), + LocalVariable(3), + Load( + pointer: 10, + ), + Literal(U32(16)), + Binary( + op: Less, + left: 11, + right: 12, + ), + GlobalVariable(0), + Load( + pointer: 1, + ), + Literal(U32(16)), + Binary( + op: Multiply, + left: 15, + right: 16, + ), + Load( + pointer: 10, + ), + Binary( + op: Add, + left: 17, + right: 18, + ), + Access( + base: 14, + index: 19, + ), + Load( + pointer: 20, + ), + Splat( + size: Quad, + value: 21, + ), + Load( + pointer: 5, + ), + Binary( + op: Add, + left: 23, + right: 22, + ), + Load( + pointer: 8, + ), + Binary( + op: Add, + left: 25, + right: 22, + ), + Literal(U32(1)), + Load( + pointer: 10, + ), + Binary( + op: Add, + left: 28, + right: 27, + ), + GlobalVariable(1), + Load( + pointer: 1, + ), + Literal(U32(2)), + Binary( + op: Multiply, + left: 31, + right: 32, + ), + Access( + base: 30, + index: 33, + ), + AccessIndex( + base: 5, + index: 0, + ), + Load( + pointer: 35, + ), + GlobalVariable(1), + Load( + pointer: 1, + ), + Literal(U32(2)), + Binary( + op: Multiply, + left: 38, + right: 39, + ), + Literal(U32(1)), + Binary( + op: Add, + left: 40, + right: 41, + ), + Access( + base: 37, + index: 42, + ), + AccessIndex( + base: 8, + index: 0, + ), + Load( + pointer: 44, + ), + Literal(U32(1)), + Load( + pointer: 1, + ), + Binary( + op: Add, + left: 47, + right: 46, + ), + ], + named_expressions: { + 22: "v", + }, + body: [ + Loop( + body: [ + Emit(( + start: 2, + end: 3, + )), + Emit(( + start: 4, + end: 5, + )), + If( + condition: 4, + accept: [], + reject: [ + Break, + ], + ), + Block([ + Store( + pointer: 5, + value: 6, + ), + Store( + pointer: 8, + value: 7, + ), + Store( + pointer: 10, + value: 9, + ), + Loop( + body: [ + Emit(( + start: 11, + end: 12, + )), + Emit(( + start: 13, + end: 14, + )), + If( + condition: 13, + accept: [], + reject: [ + Break, + ], + ), + Block([ + Emit(( + start: 15, + end: 16, + )), + Emit(( + start: 17, + end: 23, + )), + Emit(( + start: 23, + end: 25, + )), + Store( + pointer: 5, + value: 24, + ), + Emit(( + start: 25, + end: 27, + )), + Store( + pointer: 8, + value: 26, + ), + ]), + ], + continuing: [ + Emit(( + start: 28, + end: 30, + )), + Store( + pointer: 10, + value: 29, + ), + ], + break_if: None, + ), + Emit(( + start: 31, + end: 32, + )), + Emit(( + start: 33, + end: 37, + )), + Store( + pointer: 34, + value: 36, + ), + Emit(( + start: 38, + end: 39, + )), + Emit(( + start: 40, + end: 41, + )), + Emit(( + start: 42, + end: 46, + )), + Store( + pointer: 43, + value: 45, + ), + ]), + ], + continuing: [ + Emit(( + start: 47, + end: 49, + )), + Store( + pointer: 1, + value: 48, + ), + ], + break_if: None, + ), + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + mesh_info: None, + task_payload: None, + incoming_ray_payload: None, + ), + ], + diagnostic_filters: [], + diagnostic_filter_leaf: None, + doc_comments: None, +) \ No newline at end of file diff --git a/naga/tests/out/msl/wgsl-9489-loop-local-var-init.metal b/naga/tests/out/msl/wgsl-9489-loop-local-var-init.metal new file mode 100644 index 00000000000..1f320de1102 --- /dev/null +++ b/naga/tests/out/msl/wgsl-9489-loop-local-var-init.metal @@ -0,0 +1,76 @@ +// language: metal1.0 +#include +#include + +using metal::uint; + +struct type_1 { + float inner[64]; +}; +struct type_2 { + float inner[8]; +}; + +kernel void main_( + device type_1 const& input [[buffer(0)]] +, device type_2& output [[buffer(1)]] +) { + uint t = 0u; + metal::float4 acc_noinit = {}; + metal::float4 acc_init = {}; + uint d = {}; + uint2 loop_bound = uint2(4294967295u); + bool loop_init = true; + while(true) { + if (metal::all(loop_bound == uint2(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 = metal::float4 {}; + acc_init = metal::float4 {}; + d = 0u; + uint2 loop_bound_1 = uint2(4294967295u); + bool loop_init_1 = true; + while(true) { + if (metal::all(loop_bound_1 == uint2(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 = input.inner[(_e15 * 16u) + _e18]; + metal::float4 v = metal::float4(_e21); + metal::float4 _e23 = acc_noinit; + acc_noinit = _e23 + v; + metal::float4 _e25 = acc_init; + acc_init = _e25 + v; + } + } + uint _e31 = t; + float _e36 = acc_noinit.x; + output.inner[_e31 * 2u] = _e36; + uint _e38 = t; + float _e45 = acc_init.x; + output.inner[(_e38 * 2u) + 1u] = _e45; + } + } + return; +} diff --git a/naga/tests/out/spv/wgsl-9489-loop-local-var-init.spvasm b/naga/tests/out/spv/wgsl-9489-loop-local-var-init.spvasm new file mode 100644 index 00000000000..7b8a1c7dcc1 --- /dev/null +++ b/naga/tests/out/spv/wgsl-9489-loop-local-var-init.spvasm @@ -0,0 +1,175 @@ +; SPIR-V +; Version: 1.1 +; Generator: rspirv +; Bound: 116 +OpCapability Shader +OpExtension "SPV_KHR_storage_buffer_storage_class" +%1 = OpExtInstImport "GLSL.std.450" +OpMemoryModel Logical GLSL450 +OpEntryPoint GLCompute %17 "main" +OpExecutionMode %17 LocalSize 1 1 1 +OpDecorate %4 ArrayStride 4 +OpDecorate %7 ArrayStride 4 +OpDecorate %10 NonWritable +OpDecorate %10 DescriptorSet 0 +OpDecorate %10 Binding 0 +OpDecorate %11 Block +OpMemberDecorate %11 0 Offset 0 +OpDecorate %13 DescriptorSet 0 +OpDecorate %13 Binding 1 +OpDecorate %14 Block +OpMemberDecorate %14 0 Offset 0 +%2 = OpTypeVoid +%3 = OpTypeFloat 32 +%6 = OpTypeInt 32 0 +%5 = OpConstant %6 64 +%4 = OpTypeArray %3 %5 +%8 = OpConstant %6 8 +%7 = OpTypeArray %3 %8 +%9 = OpTypeVector %3 4 +%11 = OpTypeStruct %4 +%12 = OpTypePointer StorageBuffer %11 +%10 = OpVariable %12 StorageBuffer +%14 = OpTypeStruct %7 +%15 = OpTypePointer StorageBuffer %14 +%13 = OpVariable %15 StorageBuffer +%18 = OpTypeFunction %2 +%19 = OpTypePointer StorageBuffer %4 +%20 = OpConstant %6 0 +%22 = OpTypePointer StorageBuffer %7 +%24 = OpConstant %6 4 +%25 = OpConstantNull %9 +%26 = OpConstant %6 16 +%27 = OpConstant %6 1 +%28 = OpConstant %6 2 +%30 = OpTypePointer Function %6 +%32 = OpTypePointer Function %9 +%33 = OpConstantNull %9 +%35 = OpConstantNull %9 +%37 = OpConstantNull %6 +%43 = OpTypeVector %6 2 +%44 = OpTypePointer Function %43 +%45 = OpTypeBool +%46 = OpTypeVector %45 2 +%47 = OpConstantComposite %43 %20 %20 +%48 = OpConstant %6 4294967295 +%49 = OpConstantComposite %43 %48 %48 +%92 = OpTypePointer StorageBuffer %3 +%104 = OpTypePointer Function %3 +%17 = OpFunction %2 None %18 +%16 = OpLabel +%31 = OpVariable %32 Function %33 +%36 = OpVariable %30 Function %37 +%29 = OpVariable %30 Function %20 +%34 = OpVariable %32 Function %35 +%50 = OpVariable %44 Function %49 +%71 = OpVariable %44 Function %49 +%21 = OpAccessChain %19 %10 %20 +%23 = OpAccessChain %22 %13 %20 +OpBranch %38 +%38 = OpLabel +OpBranch %39 +%39 = OpLabel +OpLoopMerge %40 %42 None +OpBranch %51 +%51 = OpLabel +%52 = OpLoad %43 %50 +%53 = OpIEqual %46 %47 %52 +%54 = OpAll %45 %53 +OpSelectionMerge %55 None +OpBranchConditional %54 %40 %55 +%55 = OpLabel +%56 = OpCompositeExtract %6 %52 1 +%57 = OpIEqual %45 %56 %20 +%58 = OpSelect %6 %57 %27 %20 +%59 = OpCompositeConstruct %43 %58 %27 +%60 = OpISub %43 %52 %59 +OpStore %50 %60 +OpBranch %41 +%41 = OpLabel +%61 = OpLoad %6 %29 +%62 = OpULessThan %45 %61 %24 +OpSelectionMerge %63 None +OpBranchConditional %62 %63 %64 +%64 = OpLabel +OpBranch %40 +%63 = OpLabel +OpBranch %65 +%65 = OpLabel +OpStore %31 %25 +OpStore %34 %25 +OpStore %36 %20 +OpBranch %67 +%67 = OpLabel +OpLoopMerge %68 %70 None +OpBranch %72 +%72 = OpLabel +%73 = OpLoad %43 %71 +%74 = OpIEqual %46 %47 %73 +%75 = OpAll %45 %74 +OpSelectionMerge %76 None +OpBranchConditional %75 %68 %76 +%76 = OpLabel +%77 = OpCompositeExtract %6 %73 1 +%78 = OpIEqual %45 %77 %20 +%79 = OpSelect %6 %78 %27 %20 +%80 = OpCompositeConstruct %43 %79 %27 +%81 = OpISub %43 %73 %80 +OpStore %71 %81 +OpBranch %69 +%69 = OpLabel +%82 = OpLoad %6 %36 +%83 = OpULessThan %45 %82 %26 +OpSelectionMerge %84 None +OpBranchConditional %83 %84 %85 +%85 = OpLabel +OpBranch %68 +%84 = OpLabel +OpBranch %86 +%86 = OpLabel +%88 = OpLoad %6 %29 +%89 = OpIMul %6 %88 %26 +%90 = OpLoad %6 %36 +%91 = OpIAdd %6 %89 %90 +%93 = OpAccessChain %92 %21 %91 +%94 = OpLoad %3 %93 +%95 = OpCompositeConstruct %9 %94 %94 %94 %94 +%96 = OpLoad %9 %31 +%97 = OpFAdd %9 %96 %95 +OpStore %31 %97 +%98 = OpLoad %9 %34 +%99 = OpFAdd %9 %98 %95 +OpStore %34 %99 +OpBranch %87 +%87 = OpLabel +OpBranch %70 +%70 = OpLabel +%100 = OpLoad %6 %36 +%101 = OpIAdd %6 %100 %27 +OpStore %36 %101 +OpBranch %67 +%68 = OpLabel +%102 = OpLoad %6 %29 +%103 = OpIMul %6 %102 %28 +%105 = OpAccessChain %104 %31 %20 +%106 = OpLoad %3 %105 +%107 = OpAccessChain %92 %23 %103 +OpStore %107 %106 +%108 = OpLoad %6 %29 +%109 = OpIMul %6 %108 %28 +%110 = OpIAdd %6 %109 %27 +%111 = OpAccessChain %104 %34 %20 +%112 = OpLoad %3 %111 +%113 = OpAccessChain %92 %23 %110 +OpStore %113 %112 +OpBranch %66 +%66 = OpLabel +OpBranch %42 +%42 = OpLabel +%114 = OpLoad %6 %29 +%115 = OpIAdd %6 %114 %27 +OpStore %29 %115 +OpBranch %39 +%40 = OpLabel +OpReturn +OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/wgsl/wgsl-9489-loop-local-var-init.wgsl b/naga/tests/out/wgsl/wgsl-9489-loop-local-var-init.wgsl new file mode 100644 index 00000000000..d756815c48b --- /dev/null +++ b/naga/tests/out/wgsl/wgsl-9489-loop-local-var-init.wgsl @@ -0,0 +1,57 @@ +@group(0) @binding(0) +var input: array; +@group(0) @binding(1) +var output: array; + +@compute @workgroup_size(1, 1, 1) +fn main() { + var t: u32 = 0u; + var acc_noinit: vec4; + var acc_init: vec4; + var d: u32; + + loop { + let _e2 = t; + if (_e2 < 4u) { + } else { + break; + } + { + acc_noinit = vec4(); + acc_init = vec4(); + d = 0u; + loop { + let _e11 = d; + if (_e11 < 16u) { + } else { + break; + } + { + let _e15 = t; + let _e18 = d; + let _e21 = input[((_e15 * 16u) + _e18)]; + let v = vec4(_e21); + let _e23 = acc_noinit; + acc_noinit = (_e23 + v); + let _e25 = acc_init; + acc_init = (_e25 + v); + } + continuing { + let _e28 = d; + d = (_e28 + 1u); + } + } + let _e31 = t; + let _e36 = acc_noinit.x; + output[(_e31 * 2u)] = _e36; + let _e38 = t; + let _e45 = acc_init.x; + output[((_e38 * 2u) + 1u)] = _e45; + } + continuing { + let _e47 = t; + t = (_e47 + 1u); + } + } + return; +}