[naga] Restore return statement at end of functions without return types (#7021)

CHANGELOG.md

@@ -68,6 +68,10 @@ By @brodycj in [#6924](https://github.com/gfx-rs/wgpu/pull/6924).
 
 ### Bug Fixes
 
+#### Naga
+
+- Fix some instances of functions which have a return type but don't return a value being incorrectly validated. By @jamienicol in [#7013](https://github.com/gfx-rs/wgpu/pull/7013).
+
 #### General
 
 - Avoid overflow in query set bounds check validation. By @ErichDonGubler in [#6933](https://github.com/gfx-rs/wgpu/pull/6933).

naga/src/front/glsl/functions.rs

@@ -1031,9 +1031,7 @@ impl Frontend {
         result: Option<FunctionResult>,
         meta: Span,
     ) {
-        if result.is_some() {
-            ensure_block_returns(&mut ctx.body);
-        }
+        ensure_block_returns(&mut ctx.body);
 
         let void = result.is_none();
 

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

@@ -1307,9 +1307,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
             global_expression_kind_tracker: ctx.global_expression_kind_tracker,
         };
         let mut body = self.block(&f.body, false, &mut stmt_ctx)?;
-        if function.result.is_some() {
-            ensure_block_returns(&mut body);
-        }
+        ensure_block_returns(&mut body);
 
         function.body = body;
         function.named_expressions = named_expressions

naga/tests/out/glsl/6772-unpack-expr-accesses.main.Compute.glsl

@@ -9,5 +9,6 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 void main() {
     int phony = ivec4(bitfieldExtract(int(12u), 0, 8), bitfieldExtract(int(12u), 8, 8), bitfieldExtract(int(12u), 16, 8), bitfieldExtract(int(12u), 24, 8))[2];
     uint phony_1 = uvec4(bitfieldExtract(12u, 0, 8), bitfieldExtract(12u, 8, 8), bitfieldExtract(12u, 16, 8), bitfieldExtract(12u, 24, 8)).y;
+    return;
 }
 

naga/tests/out/glsl/access.assign_through_ptr.Compute.glsl

@@ -44,10 +44,12 @@ float test_arr_as_arg(float a[5][10]) {
 
 void assign_through_ptr_fn(inout uint p) {
     p = 42u;
+    return;
 }
 
 void assign_array_through_ptr_fn(inout vec4 foo_2[2]) {
     foo_2 = vec4[2](vec4(1.0), vec4(2.0));
+    return;
 }
 
 uint fetch_arg_ptr_member(inout AssignToMember p_1) {
@@ -57,6 +59,7 @@ uint fetch_arg_ptr_member(inout AssignToMember p_1) {
 
 void assign_to_arg_ptr_member(inout AssignToMember p_2) {
     p_2.x = 10u;
+    return;
 }
 
 uint fetch_arg_ptr_array_element(inout uint p_3[4]) {
@@ -66,6 +69,7 @@ uint fetch_arg_ptr_array_element(inout uint p_3[4]) {
 
 void assign_to_arg_ptr_array_element(inout uint p_4[4]) {
     p_4[1] = 10u;
+    return;
 }
 
 bool index_ptr(bool value) {
@@ -110,5 +114,6 @@ void main() {
     vec4 arr[2] = vec4[2](vec4(6.0), vec4(7.0));
     assign_through_ptr_fn(val);
     assign_array_through_ptr_fn(arr);
+    return;
 }
 

naga/tests/out/glsl/access.assign_to_ptr_components.Compute.glsl

@@ -44,10 +44,12 @@ float test_arr_as_arg(float a[5][10]) {
 
 void assign_through_ptr_fn(inout uint p) {
     p = 42u;
+    return;
 }
 
 void assign_array_through_ptr_fn(inout vec4 foo_2[2]) {
     foo_2 = vec4[2](vec4(1.0), vec4(2.0));
+    return;
 }
 
 uint fetch_arg_ptr_member(inout AssignToMember p_1) {
@@ -57,6 +59,7 @@ uint fetch_arg_ptr_member(inout AssignToMember p_1) {
 
 void assign_to_arg_ptr_member(inout AssignToMember p_2) {
     p_2.x = 10u;
+    return;
 }
 
 uint fetch_arg_ptr_array_element(inout uint p_3[4]) {
@@ -66,6 +69,7 @@ uint fetch_arg_ptr_array_element(inout uint p_3[4]) {
 
 void assign_to_arg_ptr_array_element(inout uint p_4[4]) {
     p_4[1] = 10u;
+    return;
 }
 
 bool index_ptr(bool value) {
@@ -112,5 +116,6 @@ void main() {
     uint _e1 = fetch_arg_ptr_member(s1_);
     assign_to_arg_ptr_array_element(a1_);
     uint _e3 = fetch_arg_ptr_array_element(a1_);
+    return;
 }
 

naga/tests/out/glsl/access.foo_frag.Fragment.glsl

@@ -54,10 +54,12 @@ float test_arr_as_arg(float a[5][10]) {
 
 void assign_through_ptr_fn(inout uint p) {
     p = 42u;
+    return;
 }
 
 void assign_array_through_ptr_fn(inout vec4 foo_2[2]) {
     foo_2 = vec4[2](vec4(1.0), vec4(2.0));
+    return;
 }
 
 uint fetch_arg_ptr_member(inout AssignToMember p_1) {
@@ -67,6 +69,7 @@ uint fetch_arg_ptr_member(inout AssignToMember p_1) {
 
 void assign_to_arg_ptr_member(inout AssignToMember p_2) {
     p_2.x = 10u;
+    return;
 }
 
 uint fetch_arg_ptr_array_element(inout uint p_3[4]) {
@@ -76,6 +79,7 @@ uint fetch_arg_ptr_array_element(inout uint p_3[4]) {
 
 void assign_to_arg_ptr_array_element(inout uint p_4[4]) {
     p_4[1] = 10u;
+    return;
 }
 
 bool index_ptr(bool value) {

naga/tests/out/glsl/access.foo_vert.Vertex.glsl

@@ -77,6 +77,7 @@ void test_matrix_within_struct_accesses() {
     int _e85 = idx;
     int _e87 = idx;
     t.m[_e85][_e87] = 40.0;
+    return;
 }
 
 void test_matrix_within_array_within_struct_accesses() {
@@ -112,6 +113,7 @@ void test_matrix_within_array_within_struct_accesses() {
     int _e100 = idx_1;
     int _e102 = idx_1;
     t_1.am[0][_e100][_e102] = 40.0;
+    return;
 }
 
 float read_from_private(inout float foo_1) {
@@ -125,10 +127,12 @@ float test_arr_as_arg(float a[5][10]) {
 
 void assign_through_ptr_fn(inout uint p) {
     p = 42u;
+    return;
 }
 
 void assign_array_through_ptr_fn(inout vec4 foo_2[2]) {
     foo_2 = vec4[2](vec4(1.0), vec4(2.0));
+    return;
 }
 
 uint fetch_arg_ptr_member(inout AssignToMember p_1) {
@@ -138,6 +142,7 @@ uint fetch_arg_ptr_member(inout AssignToMember p_1) {
 
 void assign_to_arg_ptr_member(inout AssignToMember p_2) {
     p_2.x = 10u;
+    return;
 }
 
 uint fetch_arg_ptr_array_element(inout uint p_3[4]) {
@@ -147,6 +152,7 @@ uint fetch_arg_ptr_array_element(inout uint p_3[4]) {
 
 void assign_to_arg_ptr_array_element(inout uint p_4[4]) {
     p_4[1] = 10u;
+    return;
 }
 
 bool index_ptr(bool value) {

naga/tests/out/glsl/array-in-ctor.cs_main.Compute.glsl

@@ -13,5 +13,6 @@ layout(std430) readonly buffer Ah_block_0Compute { Ah _group_0_binding_0_cs; };
 
 void main() {
     Ah ah_1 = _group_0_binding_0_cs;
+    return;
 }
 

naga/tests/out/glsl/atomicOps.cs_main.Compute.glsl

@@ -127,5 +127,6 @@ void main() {
     int _e295 = atomicExchange(workgroup_atomic_arr[1], 1);
     uint _e299 = atomicExchange(workgroup_struct.atomic_scalar, 1u);
     int _e304 = atomicExchange(workgroup_struct.atomic_arr[1], 1);
+    return;
 }
 

naga/tests/out/glsl/atomicTexture.cs_main.Compute.glsl

@@ -22,5 +22,6 @@ void main() {
     imageAtomicAnd(_group_0_binding_1_cs, ivec2(0, 0), 1);
     imageAtomicOr(_group_0_binding_1_cs, ivec2(0, 0), 1);
     imageAtomicXor(_group_0_binding_1_cs, ivec2(0, 0), 1);
+    return;
 }
 

naga/tests/out/glsl/bitcast.main.Compute.glsl

@@ -34,5 +34,6 @@ void main() {
     f3_ = intBitsToFloat(_e41);
     ivec4 _e43 = i4_;
     f4_ = intBitsToFloat(_e43);
+    return;
 }
 

naga/tests/out/glsl/bits.main.Compute.glsl

@@ -129,5 +129,6 @@ void main() {
     u3_ = bitfieldReverse(_e168);
     uvec4 _e170 = u4_;
     u4_ = bitfieldReverse(_e170);
+    return;
 }
 

naga/tests/out/glsl/boids.main.Compute.glsl

@@ -150,5 +150,6 @@ void main() {
     _group_0_binding_2_cs.particles[index].pos = _e174;
     vec2 _e179 = vVel;
     _group_0_binding_2_cs.particles[index].vel = _e179;
+    return;
 }
 

naga/tests/out/glsl/bounds-check-image-restrict.fragment_shader.Fragment.glsl

@@ -58,22 +58,27 @@ vec4 test_textureLoad_multisampled_2d(ivec2 coords_5, int _sample) {
 
 void test_textureStore_1d(int coords_10, vec4 value) {
     imageStore(_group_0_binding_8_fs, coords_10, value);
+    return;
 }
 
 void test_textureStore_2d(ivec2 coords_11, vec4 value_1) {
     imageStore(_group_0_binding_9_fs, coords_11, value_1);
+    return;
 }
 
 void test_textureStore_2d_array_u(ivec2 coords_12, uint array_index, vec4 value_2) {
     imageStore(_group_0_binding_10_fs, ivec3(coords_12, array_index), value_2);
+    return;
 }
 
 void test_textureStore_2d_array_s(ivec2 coords_13, int array_index_1, vec4 value_3) {
     imageStore(_group_0_binding_10_fs, ivec3(coords_13, array_index_1), value_3);
+    return;
 }
 
 void test_textureStore_3d(ivec3 coords_14, vec4 value_4) {
     imageStore(_group_0_binding_11_fs, coords_14, value_4);
+    return;
 }
 
 void main() {

naga/tests/out/glsl/bounds-check-image-rzsw.fragment_shader.Fragment.glsl

@@ -52,22 +52,27 @@ vec4 test_textureLoad_multisampled_2d(ivec2 coords_5, int _sample) {
 
 void test_textureStore_1d(int coords_10, vec4 value) {
     imageStore(_group_0_binding_8_fs, coords_10, value);
+    return;
 }
 
 void test_textureStore_2d(ivec2 coords_11, vec4 value_1) {
     imageStore(_group_0_binding_9_fs, coords_11, value_1);
+    return;
 }
 
 void test_textureStore_2d_array_u(ivec2 coords_12, uint array_index, vec4 value_2) {
     imageStore(_group_0_binding_10_fs, ivec3(coords_12, array_index), value_2);
+    return;
 }
 
 void test_textureStore_2d_array_s(ivec2 coords_13, int array_index_1, vec4 value_3) {
     imageStore(_group_0_binding_10_fs, ivec3(coords_13, array_index_1), value_3);
+    return;
 }
 
 void test_textureStore_3d(ivec3 coords_14, vec4 value_4) {
     imageStore(_group_0_binding_11_fs, coords_14, value_4);
+    return;
 }
 
 void main() {

naga/tests/out/glsl/break-if.main.Compute.glsl

@@ -16,6 +16,7 @@ void breakIfEmpty() {
         }
         loop_init = false;
     }
+    return;
 }
 
 void breakIfEmptyBody(bool a) {
@@ -34,6 +35,7 @@ void breakIfEmptyBody(bool a) {
         }
         loop_init_1 = false;
     }
+    return;
 }
 
 void breakIf(bool a_1) {
@@ -52,6 +54,7 @@ void breakIf(bool a_1) {
         bool _e2 = d;
         e = (a_1 != _e2);
     }
+    return;
 }
 
 void breakIfSeparateVariable() {
@@ -68,8 +71,10 @@ void breakIfSeparateVariable() {
         uint _e3 = counter;
         counter = (_e3 + 1u);
     }
+    return;
 }
 
 void main() {
+    return;
 }
 

naga/tests/out/glsl/const-exprs.main.Compute.glsl

@@ -23,14 +23,17 @@ const bvec2 compare_vec = bvec2(true, false);
 
 void swizzle_of_compose() {
     ivec4 out_ = ivec4(4, 3, 2, 1);
+    return;
 }
 
 void index_of_compose() {
     int out_1 = 2;
+    return;
 }
 
 void compose_three_deep() {
     int out_2 = 6;
+    return;
 }
 
 void non_constant_initializers() {
@@ -48,18 +51,22 @@ void non_constant_initializers() {
     int _e10 = y;
     int _e11 = z;
     out_3 = ivec4(_e8, _e9, _e10, _e11);
+    return;
 }
 
 void splat_of_constant() {
     ivec4 out_4 = ivec4(-4, -4, -4, -4);
+    return;
 }
 
 void compose_of_constant() {
     ivec4 out_5 = ivec4(-4, -4, -4, -4);
+    return;
 }
 
 void compose_of_splat() {
     vec4 x_1 = vec4(2.0, 1.0, 1.0, 1.0);
+    return;
 }
 
 uint map_texture_kind(int texture_kind) {
@@ -87,5 +94,6 @@ void main() {
     splat_of_constant();
     compose_of_constant();
     compose_of_splat();
+    return;
 }
 

naga/tests/out/glsl/constructors.main.Compute.glsl

@@ -33,5 +33,6 @@ void main() {
     int cit2_[4] = int[4](0, 1, 2, 3);
     uvec2 ic4_ = uvec2(0u, 0u);
     mat2x3 ic5_ = mat2x3(vec3(0.0, 0.0, 0.0), vec3(0.0, 0.0, 0.0));
+    return;
 }
 

naga/tests/out/glsl/control-flow.main.Compute.glsl

@@ -35,6 +35,7 @@ void loop_switch_continue(int x) {
             }
         }
     }
+    return;
 }
 
 void loop_switch_continue_nesting(int x_1, int y, int z) {
@@ -92,6 +93,7 @@ void loop_switch_continue_nesting(int x_1, int y, int z) {
             continue;
         }
     }
+    return;
 }
 
 void loop_switch_omit_continue_variable_checks(int x_2, int y_1, int z_1, int w) {
@@ -137,6 +139,7 @@ void loop_switch_omit_continue_variable_checks(int x_2, int y_1, int z_1, int w)
             }
         }
     }
+    return;
 }
 
 void main() {
@@ -190,18 +193,18 @@ void main() {
         }
         case 2: {
             pos = 1;
-            break;
+            return;
         }
         case 3: {
             pos = 2;
-            break;
+            return;
         }
         case 4: {
-            break;
+            return;
         }
         default: {
             pos = 3;
-            break;
+            return;
         }
     }
 }

naga/tests/out/glsl/cross.main.Compute.glsl

@@ -8,5 +8,6 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 
 void main() {
     vec3 a = cross(vec3(0.0, 1.0, 2.0), vec3(0.0, 1.0, 2.0));
+    return;
 }
 

naga/tests/out/glsl/empty.main.Compute.glsl

@@ -7,5 +7,6 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 
 
 void main() {
+    return;
 }
 

naga/tests/out/glsl/f64.main.Compute.glsl

@@ -14,5 +14,6 @@ double f(double x) {
 
 void main() {
     double _e1 = f(6.0LF);
+    return;
 }
 

naga/tests/out/glsl/functions-webgl.main.Fragment.glsl

@@ -13,5 +13,6 @@ vec2 test_fma() {
 
 void main() {
     vec2 _e0 = test_fma();
+    return;
 }
 

naga/tests/out/glsl/functions.main.Compute.glsl

@@ -29,5 +29,6 @@ int test_integer_dot_product() {
 void main() {
     vec2 _e0 = test_fma();
     int _e1 = test_integer_dot_product();
+    return;
 }
 

naga/tests/out/glsl/globals.main.Compute.glsl

@@ -31,6 +31,7 @@ uniform type_15_block_6Compute { mat4x2 _group_0_binding_7_cs[2][2]; };
 
 
 void test_msl_packed_vec3_as_arg(vec3 arg) {
+    return;
 }
 
 void test_msl_packed_vec3_() {
@@ -48,6 +49,7 @@ void test_msl_packed_vec3_() {
     vec3 mvm1_ = (mat3x3(0.0) * data.v3_);
     vec3 svm0_ = (data.v3_ * 2.0);
     vec3 svm1_ = (2.0 * data.v3_);
+    return;
 }
 
 void main() {
@@ -77,5 +79,6 @@ void main() {
     _group_0_binding_1_cs.v1_ = 4.0;
     wg[1] = float(uint(_group_0_binding_2_cs.length()));
     at_1 = 2u;
+    return;
 }
 

naga/tests/out/glsl/image.main.Compute.glsl

@@ -33,5 +33,6 @@ void main() {
     uvec4 value7u = texelFetch(_group_0_binding_7_cs, int(uint(local_id.x)), int(local_id.z));
     imageStore(_group_0_binding_2_cs, itc.x, ((((value1_ + value2_) + value4_) + value5_) + value6_));
     imageStore(_group_0_binding_2_cs, int(uint(itc.x)), ((((value1u + value2u) + value4u) + value5u) + value6u));
+    return;
 }
 

naga/tests/out/glsl/interpolate_compat.frag_main.Fragment.glsl

@@ -25,5 +25,6 @@ smooth in float _vs2fs_location11;
 
 void main() {
     FragmentInput val = FragmentInput(gl_FragCoord, _vs2fs_location0, _vs2fs_location2, _vs2fs_location3, _vs2fs_location4, _vs2fs_location6, _vs2fs_location7, _vs2fs_location8, _vs2fs_location9, _vs2fs_location10, _vs2fs_location11);
+    return;
 }
 

naga/tests/out/glsl/invariant.fs.Fragment.glsl

@@ -6,5 +6,6 @@ precision highp int;
 
 void main() {
     vec4 position = gl_FragCoord;
+    return;
 }
 

naga/tests/out/glsl/math-functions.main.Fragment.glsl

@@ -94,5 +94,6 @@ void main() {
     vec3 quantizeToF16_c = vec3(unpackHalf2x16(packHalf2x16(_e125.xy)), unpackHalf2x16(packHalf2x16(_e125.zz)).x);
     vec4 _e131 = vec4(1.0, 1.0, 1.0, 1.0);
     vec4 quantizeToF16_d = vec4(unpackHalf2x16(packHalf2x16(_e131.xy)), unpackHalf2x16(packHalf2x16(_e131.zw)));
+    return;
 }
 

naga/tests/out/glsl/multiview.main.Fragment.glsl

@@ -7,5 +7,6 @@ precision highp int;
 
 void main() {
     int view_index = gl_ViewIndex;
+    return;
 }
 

naga/tests/out/glsl/multiview_webgl.main.Fragment.glsl

@@ -7,5 +7,6 @@ precision highp int;
 
 void main() {
     int view_index = int(gl_ViewID_OVR);
+    return;
 }
 

naga/tests/out/glsl/operators.main.Compute.glsl

@@ -55,6 +55,7 @@ void logical() {
     bvec3 bitwise_or1_ = bvec3(bvec3(true).x || bvec3(false).x, bvec3(true).y || bvec3(false).y, bvec3(true).z || bvec3(false).z);
     bool bitwise_and0_ = (true && false);
     bvec4 bitwise_and1_ = bvec4(bvec4(true).x && bvec4(false).x, bvec4(true).y && bvec4(false).y, bvec4(true).z && bvec4(false).z, bvec4(true).w && bvec4(false).w);
+    return;
 }
 
 void arithmetic() {
@@ -130,6 +131,7 @@ void arithmetic() {
     vec3 mul_vector0_ = (mat4x3(0.0) * vec4(1.0));
     vec4 mul_vector1_ = (vec3(2.0) * mat4x3(0.0));
     mat3x3 mul = (mat4x3(0.0) * mat3x4(0.0));
+    return;
 }
 
 void bit() {
@@ -157,6 +159,7 @@ void bit() {
     uint shr1_ = (2u >> 1u);
     ivec2 shr2_ = (ivec2(2) >> uvec2(1u));
     uvec3 shr3_ = (uvec3(2u) >> uvec3(1u));
+    return;
 }
 
 void comparison() {
@@ -196,6 +199,7 @@ void comparison() {
     bvec2 gte3_ = greaterThanEqual(ivec2(2), ivec2(1));
     bvec3 gte4_ = greaterThanEqual(uvec3(2u), uvec3(1u));
     bvec4 gte5_ = greaterThanEqual(vec4(2.0), vec4(1.0));
+    return;
 }
 
 void assignment() {
@@ -232,6 +236,7 @@ void assignment() {
     vec0_[1] = (_e37 + 1);
     int _e41 = vec0_[1];
     vec0_[1] = (_e41 - 1);
+    return;
 }
 
 void negation_avoids_prefix_decrement() {
@@ -243,6 +248,7 @@ void negation_avoids_prefix_decrement() {
     int p5_ = -(-(-(-(1))));
     int p6_ = -(-(-(-(-(1)))));
     int p7_ = -(-(-(-(-(1)))));
+    return;
 }
 
 void main() {
@@ -255,5 +261,6 @@ void main() {
     bit();
     comparison();
     assignment();
+    return;
 }
 

naga/tests/out/glsl/overrides.main.Compute.glsl

@@ -27,5 +27,6 @@ void main() {
     float _e9 = gain_x_10_;
     gain_x_100_ = (_e9 * 10.0);
     store_override = gain;
+    return;
 }
 

naga/tests/out/glsl/phony_assignment.main.Compute.glsl

@@ -19,5 +19,6 @@ void main() {
     int _e6 = five();
     int _e7 = five();
     float phony_2 = _group_0_binding_0_cs;
+    return;
 }
 

naga/tests/out/glsl/separate-entry-points.compute.Compute.glsl

@@ -11,9 +11,11 @@ void barriers() {
     barrier();
     memoryBarrierShared();
     barrier();
+    return;
 }
 
 void main() {
     barriers();
+    return;
 }
 

naga/tests/out/glsl/separate-entry-points.fragment.Fragment.glsl

@@ -9,6 +9,7 @@ void derivatives() {
     float x = dFdx(0.0);
     float y = dFdy(0.0);
     float width = fwidth(0.0);
+    return;
 }
 
 void main() {

naga/tests/out/glsl/struct-layout.needs_padding_comp.Compute.glsl

@@ -25,5 +25,6 @@ void main() {
     x_1 = _e2;
     NeedsPadding _e4 = _group_0_binding_3_cs;
     x_1 = _e4;
+    return;
 }
 

naga/tests/out/glsl/struct-layout.no_padding_comp.Compute.glsl

@@ -25,5 +25,6 @@ void main() {
     x = _e2;
     NoPadding _e4 = _group_0_binding_1_cs;
     x = _e4;
+    return;
 }
 

naga/tests/out/glsl/subgroup-operations.main.Compute.glsl

@@ -40,5 +40,6 @@ void main() {
     uint _e35 = subgroupShuffleDown(subgroup_invocation_id, 1u);
     uint _e37 = subgroupShuffleUp(subgroup_invocation_id, 1u);
     uint _e41 = subgroupShuffleXor(subgroup_invocation_id, (sizes.subgroup_size - 1u));
+    return;
 }
 

naga/tests/out/glsl/variations.main.Fragment.glsl

@@ -10,6 +10,7 @@ void main_1() {
     ivec2 sizeCube = ivec2(0);
     float a = 1.0;
     sizeCube = ivec2(uvec2(textureSize(_group_0_binding_0_fs, 0).xy));
+    return;
 }
 
 void main() {

naga/tests/out/glsl/workgroup-uniform-load.test_workgroupUniformLoad.Compute.glsl

@@ -25,6 +25,9 @@ void main() {
     if ((_e4 > 10)) {
         memoryBarrierShared();
         barrier();
+        return;
+    } else {
+        return;
     }
 }
 

naga/tests/out/glsl/workgroup-var-init.main.Compute.glsl

@@ -23,5 +23,6 @@ void main() {
     barrier();
     uint _e3[512] = w_mem.arr;
     _group_0_binding_0_cs = _e3;
+    return;
 }
 

naga/tests/out/hlsl/6772-unpack-expr-accesses.hlsl

@@ -3,4 +3,5 @@ void main()
 {
     int phony = (int4(12u, 12u >> 8, 12u >> 16, 12u >> 24) << 24 >> 24)[2];
     uint phony_1 = (uint4(12u, 12u >> 8, 12u >> 16, 12u >> 24) << 24 >> 24).y;
+    return;
 }

naga/tests/out/hlsl/access.hlsl

@@ -166,6 +166,7 @@ void test_matrix_within_struct_accesses()
     int _e85 = idx;
     int _e87 = idx;
     SetMatScalarmOnBaz(t, 40.0, _e85, _e87);
+    return;
 }
 
 MatCx2InArray ConstructMatCx2InArray(float4x2 arg0[2]) {
@@ -214,6 +215,7 @@ void test_matrix_within_array_within_struct_accesses()
     int _e100 = idx_1;
     int _e102 = idx_1;
     __set_el_of_mat4x2(t_1.am[0], _e100, _e102, 40.0);
+    return;
 }
 
 float read_from_private(inout float foo_1)
@@ -230,6 +232,7 @@ float test_arr_as_arg(float a[5][10])
 void assign_through_ptr_fn(inout uint p)
 {
     p = 42u;
+    return;
 }
 
 typedef float4 ret_Constructarray2_float4_[2];
@@ -241,6 +244,7 @@ ret_Constructarray2_float4_ Constructarray2_float4_(float4 arg0, float4 arg1) {
 void assign_array_through_ptr_fn(inout float4 foo_2[2])
 {
     foo_2 = Constructarray2_float4_((1.0).xxxx, (2.0).xxxx);
+    return;
 }
 
 uint fetch_arg_ptr_member(inout AssignToMember p_1)
@@ -252,6 +256,7 @@ uint fetch_arg_ptr_member(inout AssignToMember p_1)
 void assign_to_arg_ptr_member(inout AssignToMember p_2)
 {
     p_2.x = 10u;
+    return;
 }
 
 uint fetch_arg_ptr_array_element(inout uint p_3[4])
@@ -263,6 +268,7 @@ uint fetch_arg_ptr_array_element(inout uint p_3[4])
 void assign_to_arg_ptr_array_element(inout uint p_4[4])
 {
     p_4[1] = 10u;
+    return;
 }
 
 typedef bool ret_Constructarray1_bool_[1];
@@ -403,6 +409,7 @@ void assign_through_ptr()
 
     assign_through_ptr_fn(val);
     assign_array_through_ptr_fn(arr);
+    return;
 }
 
 [numthreads(1, 1, 1)]
@@ -415,4 +422,5 @@ void assign_to_ptr_components()
     const uint _e1 = fetch_arg_ptr_member(s1_);
     assign_to_arg_ptr_array_element(a1_);
     const uint _e3 = fetch_arg_ptr_array_element(a1_);
+    return;
 }

naga/tests/out/hlsl/array-in-ctor.hlsl

@@ -20,4 +20,5 @@ Ah ConstructAh(float arg0[2]) {
 void cs_main()
 {
     Ah ah_1 = ConstructAh(Constructarray2_float_(asfloat(ah.Load(0+0)), asfloat(ah.Load(0+4))));
+    return;
 }

naga/tests/out/hlsl/atomicOps-int64-min-max.hlsl

@@ -31,4 +31,5 @@ void cs_main(uint3 id : SV_GroupThreadID)
     storage_atomic_arr.InterlockedMin64(8, (1uL + _e24));
     storage_struct.InterlockedMin64(0, 1uL);
     storage_struct.InterlockedMin64(8+8, uint64_t(id.x));
+    return;
 }

naga/tests/out/hlsl/atomicOps-int64.hlsl

@@ -114,4 +114,5 @@ void cs_main(uint3 id : SV_GroupThreadID, uint3 __local_invocation_id : SV_Group
     int64_t _e279; InterlockedExchange(workgroup_atomic_arr[1], 1L, _e279);
     uint64_t _e283; InterlockedExchange(workgroup_struct.atomic_scalar, 1uL, _e283);
     int64_t _e288; InterlockedExchange(workgroup_struct.atomic_arr[1], 1L, _e288);
+    return;
 }

naga/tests/out/hlsl/atomicOps.hlsl

@@ -107,4 +107,5 @@ void cs_main(uint3 id : SV_GroupThreadID, uint3 __local_invocation_id : SV_Group
     int _e295; InterlockedExchange(workgroup_atomic_arr[1], 1, _e295);
     uint _e299; InterlockedExchange(workgroup_struct.atomic_scalar, 1u, _e299);
     int _e304; InterlockedExchange(workgroup_struct.atomic_arr[1], 1, _e304);
+    return;
 }

naga/tests/out/hlsl/atomicTexture-int64.hlsl

@@ -13,4 +13,5 @@ void cs_main(uint3 id : SV_GroupThreadID)
     InterlockedMax(image[int2(0, 0)],1uL);
     GroupMemoryBarrierWithGroupSync();
     InterlockedMin(image[int2(0, 0)],1uL);
+    return;
 }

naga/tests/out/hlsl/atomicTexture.hlsl

@@ -23,4 +23,5 @@ void cs_main(uint3 id : SV_GroupThreadID)
     InterlockedAnd(image_s[int2(0, 0)],1);
     InterlockedOr(image_s[int2(0, 0)],1);
     InterlockedXor(image_s[int2(0, 0)],1);
+    return;
 }

naga/tests/out/hlsl/bitcast.hlsl

@@ -29,4 +29,5 @@ void main()
     f3_ = asfloat(_e41);
     int4 _e43 = i4_;
     f4_ = asfloat(_e43);
+    return;
 }

naga/tests/out/hlsl/bits.hlsl

@@ -300,4 +300,5 @@ void main()
     u3_ = reversebits(_e168);
     uint4 _e170 = u4_;
     u4_ = reversebits(_e170);
+    return;
 }

naga/tests/out/hlsl/boids.hlsl

@@ -140,4 +140,5 @@ void main(uint3 global_invocation_id : SV_DispatchThreadID)
     particlesDst.Store2(0+index*16+0, asuint(_e174));
     float2 _e179 = vVel;
     particlesDst.Store2(8+index*16+0, asuint(_e179));
+    return;
 }

naga/tests/out/hlsl/bounds-check-dynamic-buffer.hlsl

@@ -35,4 +35,5 @@ void main()
     out_.Store(8, asuint(_e19));
     uint _e25 = asuint(in_data_storage_g1_b0_.Load(0+i*16+__dynamic_buffer_offsets1._0));
     out_.Store(12, asuint(_e25));
+    return;
 }

naga/tests/out/hlsl/break-if.hlsl

@@ -9,6 +9,7 @@ void breakIfEmpty()
         }
         loop_init = false;
     }
+    return;
 }
 
 void breakIfEmptyBody(bool a)
@@ -29,6 +30,7 @@ void breakIfEmptyBody(bool a)
         }
         loop_init_1 = false;
     }
+    return;
 }
 
 void breakIf(bool a_1)
@@ -49,6 +51,7 @@ void breakIf(bool a_1)
         bool _e2 = d;
         e = (a_1 != _e2);
     }
+    return;
 }
 
 void breakIfSeparateVariable()
@@ -67,9 +70,11 @@ void breakIfSeparateVariable()
         uint _e3 = counter;
         counter = (_e3 + 1u);
     }
+    return;
 }
 
 [numthreads(1, 1, 1)]
 void main()
 {
+    return;
 }

naga/tests/out/hlsl/collatz.hlsl

@@ -35,4 +35,5 @@ void main(uint3 global_id : SV_DispatchThreadID)
     uint _e9 = asuint(v_indices.Load(global_id.x*4+0));
     const uint _e10 = collatz_iterations(_e9);
     v_indices.Store(global_id.x*4+0, asuint(_e10));
+    return;
 }

naga/tests/out/hlsl/const-exprs.hlsl

@@ -17,18 +17,21 @@ void swizzle_of_compose()
 {
     int4 out_ = int4(4, 3, 2, 1);
 
+    return;
 }
 
 void index_of_compose()
 {
     int out_1 = 2;
 
+    return;
 }
 
 void compose_three_deep()
 {
     int out_2 = 6;
 
+    return;
 }
 
 void non_constant_initializers()
@@ -48,24 +51,28 @@ void non_constant_initializers()
     int _e10 = y;
     int _e11 = z;
     out_3 = int4(_e8, _e9, _e10, _e11);
+    return;
 }
 
 void splat_of_constant()
 {
     int4 out_4 = int4(-4, -4, -4, -4);
 
+    return;
 }
 
 void compose_of_constant()
 {
     int4 out_5 = int4(-4, -4, -4, -4);
 
+    return;
 }
 
 void compose_of_splat()
 {
     float4 x_1 = float4(2.0, 1.0, 1.0, 1.0);
 
+    return;
 }
 
 uint map_texture_kind(int texture_kind)
@@ -96,4 +103,5 @@ void main()
     splat_of_constant();
     compose_of_constant();
     compose_of_splat();
+    return;
 }

naga/tests/out/hlsl/constructors.hlsl

@@ -88,4 +88,5 @@ void main()
     int cit2_[4] = Constructarray4_int_(0, 1, 2, 3);
     uint2 ic4_ = uint2(0u, 0u);
     float2x3 ic5_ = float2x3(float3(0.0, 0.0, 0.0), float3(0.0, 0.0, 0.0));
+    return;
 }

naga/tests/out/hlsl/control-flow.hlsl

@@ -35,6 +35,7 @@ void loop_switch_continue(int x)
             continue;
         }
     }
+    return;
 }
 
 void loop_switch_continue_nesting(int x_1, int y, int z)
@@ -107,6 +108,7 @@ void loop_switch_continue_nesting(int x_1, int y, int z)
             continue;
         }
     }
+    return;
 }
 
 void loop_switch_omit_continue_variable_checks(int x_2, int y_1, int z_1, int w)
@@ -163,6 +165,7 @@ void loop_switch_omit_continue_variable_checks(int x_2, int y_1, int z_1, int w)
             continue;
         }
     }
+    return;
 }
 
 [numthreads(1, 1, 1)]
@@ -216,18 +219,18 @@ void main(uint3 global_id : SV_DispatchThreadID)
         }
         case 2: {
             pos = 1;
-            break;
+            return;
         }
         case 3: {
             pos = 2;
-            break;
+            return;
         }
         case 4: {
-            break;
+            return;
         }
         default: {
             pos = 3;
-            break;
+            return;
         }
     }
 }

naga/tests/out/hlsl/cross.hlsl

@@ -2,4 +2,5 @@
 void main()
 {
     float3 a = cross(float3(0.0, 1.0, 2.0), float3(0.0, 1.0, 2.0));
+    return;
 }

naga/tests/out/hlsl/empty.hlsl

@@ -1,4 +1,5 @@
 [numthreads(1, 1, 1)]
 void main()
 {
+    return;
 }

naga/tests/out/hlsl/f64.hlsl

@@ -15,4 +15,5 @@ double f(double x)
 void main()
 {
     const double _e1 = f(6.0L);
+    return;
 }

naga/tests/out/hlsl/functions.hlsl

@@ -23,4 +23,5 @@ void main()
 {
     const float2 _e0 = test_fma();
     const int _e1 = test_integer_dot_product();
+    return;
 }

naga/tests/out/hlsl/globals.hlsl

@@ -68,6 +68,7 @@ cbuffer global_nested_arrays_of_matrices_4x2_ : register(b7) { __mat4x2 global_n
 
 void test_msl_packed_vec3_as_arg(float3 arg)
 {
+    return;
 }
 
 float3x3 ZeroValuefloat3x3() {
@@ -98,6 +99,7 @@ void test_msl_packed_vec3_()
     float3 mvm1_ = mul(data.v3_, ZeroValuefloat3x3());
     float3 svm0_ = (data.v3_ * 2.0);
     float3 svm1_ = (2.0 * data.v3_);
+    return;
 }
 
 uint NagaBufferLength(ByteAddressBuffer buffer)
@@ -136,4 +138,5 @@ void main(uint3 __local_invocation_id : SV_GroupThreadID)
     alignment.Store(12, asuint(4.0));
     wg[1] = float(((NagaBufferLength(dummy) - 0) / 8));
     at_1 = 2u;
+    return;
 }

naga/tests/out/hlsl/image.hlsl

@@ -50,6 +50,7 @@ void main(uint3 local_id : SV_GroupThreadID)
     uint4 value7u = image_1d_src.Load(int2(uint(local_id.x), int(local_id.z)));
     image_dst[itc.x] = ((((value1_ + value2_) + value4_) + value5_) + value6_);
     image_dst[uint(itc.x)] = ((((value1u + value2u) + value4u) + value5u) + value6u);
+    return;
 }
 
 [numthreads(16, 1, 1)]

naga/tests/out/hlsl/int64.hlsl

@@ -230,4 +230,5 @@ void main()
     const uint64_t _e3 = uint64_function(67uL);
     const int64_t _e5 = int64_function(60L);
     output.Store(224, (_e3 + _e5));
+    return;
 }

naga/tests/out/hlsl/interface.hlsl

@@ -82,6 +82,7 @@ void compute(uint3 global_id : SV_DispatchThreadID, uint3 local_id : SV_GroupThr
     }
     GroupMemoryBarrierWithGroupSync();
     output[0] = ((((global_id.x + local_id.x) + local_index) + wg_id.x) + uint3(_NagaConstants.first_vertex, _NagaConstants.first_instance, _NagaConstants.other).x);
+    return;
 }
 
 precise float4 vertex_two_structs(Input1_ in1_, Input2_ in2_) : SV_Position

naga/tests/out/hlsl/interpolate.hlsl

@@ -68,4 +68,5 @@ VertexOutput_vert_main vert_main()
 void frag_main(FragmentInput_frag_main fragmentinput_frag_main)
 {
     FragmentInput val = { fragmentinput_frag_main.position_1, fragmentinput_frag_main._flat_1, fragmentinput_frag_main.flat_first_1, fragmentinput_frag_main.flat_either_1, fragmentinput_frag_main._linear_1, fragmentinput_frag_main.linear_centroid_1, fragmentinput_frag_main.linear_sample_1, fragmentinput_frag_main.linear_center_1, fragmentinput_frag_main.perspective_1, fragmentinput_frag_main.perspective_centroid_1, fragmentinput_frag_main.perspective_sample_1, fragmentinput_frag_main.perspective_center_1 };
+    return;
 }

naga/tests/out/hlsl/interpolate_compat.hlsl

@@ -64,4 +64,5 @@ VertexOutput_vert_main vert_main()
 void frag_main(FragmentInput_frag_main fragmentinput_frag_main)
 {
     FragmentInput val = { fragmentinput_frag_main.position_1, fragmentinput_frag_main._flat_1, fragmentinput_frag_main.flat_either_1, fragmentinput_frag_main._linear_1, fragmentinput_frag_main.linear_centroid_1, fragmentinput_frag_main.linear_sample_1, fragmentinput_frag_main.linear_center_1, fragmentinput_frag_main.perspective_1, fragmentinput_frag_main.perspective_centroid_1, fragmentinput_frag_main.perspective_sample_1, fragmentinput_frag_main.perspective_center_1 };
+    return;
 }

naga/tests/out/hlsl/math-functions.hlsl

@@ -105,4 +105,5 @@ void main()
     float2 quantizeToF16_b = f16tof32(f32tof16(float2(1.0, 1.0)));
     float3 quantizeToF16_c = f16tof32(f32tof16(float3(1.0, 1.0, 1.0)));
     float4 quantizeToF16_d = f16tof32(f32tof16(float4(1.0, 1.0, 1.0, 1.0)));
+    return;
 }

naga/tests/out/hlsl/operators.hlsl

@@ -53,6 +53,7 @@ void logical()
     bool3 bitwise_or1_ = ((true).xxx | (false).xxx);
     bool bitwise_and0_ = (true & false);
     bool4 bitwise_and1_ = ((true).xxxx & (false).xxxx);
+    return;
 }
 
 float3x3 ZeroValuefloat3x3() {
@@ -141,6 +142,7 @@ void arithmetic()
     float3 mul_vector0_ = mul((1.0).xxxx, ZeroValuefloat4x3());
     float4 mul_vector1_ = mul(ZeroValuefloat4x3(), (2.0).xxx);
     float3x3 mul_ = mul(ZeroValuefloat3x4(), ZeroValuefloat4x3());
+    return;
 }
 
 void bit()
@@ -169,6 +171,7 @@ void bit()
     uint shr1_ = (2u >> 1u);
     int2 shr2_ = ((2).xx >> (1u).xx);
     uint3 shr3_ = ((2u).xxx >> (1u).xxx);
+    return;
 }
 
 void comparison()
@@ -209,6 +212,7 @@ void comparison()
     bool2 gte3_ = ((2).xx >= (1).xx);
     bool3 gte4_ = ((2u).xxx >= (1u).xxx);
     bool4 gte5_ = ((2.0).xxxx >= (1.0).xxxx);
+    return;
 }
 
 int3 ZeroValueint3() {
@@ -251,6 +255,7 @@ void assignment()
     vec0_[1] = (_e37 + 1);
     int _e41 = vec0_[1];
     vec0_[1] = (_e41 - 1);
+    return;
 }
 
 void negation_avoids_prefix_decrement()
@@ -263,6 +268,7 @@ void negation_avoids_prefix_decrement()
     int p5_ = -(-(-(-(1))));
     int p6_ = -(-(-(-(-(1)))));
     int p7_ = -(-(-(-(-(1)))));
+    return;
 }
 
 [numthreads(1, 1, 1)]
@@ -276,4 +282,5 @@ void main(uint3 id : SV_GroupID)
     bit();
     comparison();
     assignment();
+    return;
 }

naga/tests/out/hlsl/overrides.hlsl

@@ -21,4 +21,5 @@ void main()
     float _e9 = gain_x_10_;
     gain_x_100_ = (_e9 * 10.0);
     store_override = gain;
+    return;
 }

naga/tests/out/hlsl/phony_assignment.hlsl

@@ -13,4 +13,5 @@ void main(uint3 id : SV_DispatchThreadID)
     const int _e6 = five();
     const int _e7 = five();
     float phony_2 = binding;
+    return;
 }

naga/tests/out/hlsl/ray-query.hlsl

@@ -114,6 +114,7 @@ void main()
     output.Store(0, asuint(uint((_e7.kind == 0u))));
     const float3 _e18 = get_torus_normal((dir_1 * _e7.t), _e7);
     output.Store3(16, asuint(_e18));
+    return;
 }
 
 RayIntersection GetCandidateIntersection(RayQuery<RAY_FLAG_NONE> rq) {
@@ -147,4 +148,5 @@ void main_candidate()
     rq.TraceRayInline(acc_struct, ConstructRayDesc_(4u, 255u, 0.1, 100.0, pos_2, dir_2).flags, ConstructRayDesc_(4u, 255u, 0.1, 100.0, pos_2, dir_2).cull_mask, RayDescFromRayDesc_(ConstructRayDesc_(4u, 255u, 0.1, 100.0, pos_2, dir_2)));
     RayIntersection intersection_1 = GetCandidateIntersection(rq);
     output.Store(0, asuint(uint((intersection_1.kind == 3u))));
+    return;
 }

naga/tests/out/hlsl/struct-layout.hlsl

@@ -52,6 +52,7 @@ void no_padding_comp()
     x = _e2;
     NoPadding _e4 = ConstructNoPadding(asfloat(no_padding_storage.Load3(0)), asfloat(no_padding_storage.Load(12)));
     x = _e4;
+    return;
 }
 
 float4 needs_padding_frag(FragmentInput_needs_padding_frag fragmentinput_needs_padding_frag) : SV_Target0
@@ -82,4 +83,5 @@ void needs_padding_comp()
     x_1 = _e2;
     NeedsPadding _e4 = ConstructNeedsPadding(asfloat(needs_padding_storage.Load(0)), asfloat(needs_padding_storage.Load3(16)), asfloat(needs_padding_storage.Load(28)));
     x_1 = _e4;
+    return;
 }

naga/tests/out/hlsl/subgroup-operations.hlsl

@@ -34,4 +34,5 @@ void main(ComputeInput_main computeinput_main)
     const uint _e35 = WaveReadLaneAt(subgroup_invocation_id, WaveGetLaneIndex() + 1u);
     const uint _e37 = WaveReadLaneAt(subgroup_invocation_id, WaveGetLaneIndex() - 1u);
     const uint _e41 = WaveReadLaneAt(subgroup_invocation_id, WaveGetLaneIndex() ^ (sizes.subgroup_size - 1u));
+    return;
 }

naga/tests/out/hlsl/workgroup-uniform-load.hlsl

@@ -14,5 +14,8 @@ void test_workgroupUniformLoad(uint3 workgroup_id : SV_GroupID, uint3 __local_in
     GroupMemoryBarrierWithGroupSync();
     if ((_e4 > 10)) {
         GroupMemoryBarrierWithGroupSync();
+        return;
+    } else {
+        return;
     }
 }

naga/tests/out/hlsl/workgroup-var-init.hlsl

@@ -530,4 +530,5 @@ void main(uint3 __local_invocation_id : SV_GroupThreadID)
         output.Store(2040, asuint(_value2[510]));
         output.Store(2044, asuint(_value2[511]));
     }
+    return;
 }

naga/tests/out/ir/access.compact.ron

@@ -982,6 +982,9 @@
                     pointer: 88,
                     value: 89,
                 ),
+                Return(
+                    value: None,
+                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -1584,6 +1587,9 @@
                     pointer: 103,
                     value: 104,
                 ),
+                Return(
+                    value: None,
+                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -1687,6 +1693,9 @@
                     pointer: 0,
                     value: 1,
                 ),
+                Return(
+                    value: None,
+                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -1745,6 +1754,9 @@
                     pointer: 0,
                     value: 5,
                 ),
+                Return(
+                    value: None,
+                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -1817,6 +1829,9 @@
                     pointer: 1,
                     value: 2,
                 ),
+                Return(
+                    value: None,
+                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -1889,6 +1904,9 @@
                     pointer: 1,
                     value: 2,
                 ),
+                Return(
+                    value: None,
+                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -2761,6 +2779,9 @@
                         ],
                         result: None,
                     ),
+                    Return(
+                        value: None,
+                    ),
                 ],
                 diagnostic_filter_leaf: None,
             ),
@@ -2823,6 +2844,9 @@
                         ],
                         result: Some(3),
                     ),
+                    Return(
+                        value: None,
+                    ),
                 ],
                 diagnostic_filter_leaf: None,
             ),

naga/tests/out/ir/access.ron

@@ -982,6 +982,9 @@
                     pointer: 88,
                     value: 89,
                 ),
+                Return(
+                    value: None,
+                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -1584,6 +1587,9 @@
                     pointer: 103,
                     value: 104,
                 ),
+                Return(
+                    value: None,
+                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -1687,6 +1693,9 @@
                     pointer: 0,
                     value: 1,
                 ),
+                Return(
+                    value: None,
+                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -1745,6 +1754,9 @@
                     pointer: 0,
                     value: 5,
                 ),
+                Return(
+                    value: None,
+                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -1817,6 +1829,9 @@
                     pointer: 1,
                     value: 2,
                 ),
+                Return(
+                    value: None,
+                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -1889,6 +1904,9 @@
                     pointer: 1,
                     value: 2,
                 ),
+                Return(
+                    value: None,
+                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -2761,6 +2779,9 @@
                         ],
                         result: None,
                     ),
+                    Return(
+                        value: None,
+                    ),
                 ],
                 diagnostic_filter_leaf: None,
             ),
@@ -2823,6 +2844,9 @@
                         ],
                         result: Some(3),
                     ),
+                    Return(
+                        value: None,
+                    ),
                 ],
                 diagnostic_filter_leaf: None,
             ),

naga/tests/out/ir/collatz.compact.ron

@@ -325,6 +325,9 @@
                         pointer: 4,
                         value: 10,
                     ),
+                    Return(
+                        value: None,
+                    ),
                 ],
                 diagnostic_filter_leaf: None,
             ),

naga/tests/out/ir/collatz.ron

@@ -325,6 +325,9 @@
                         pointer: 4,
                         value: 10,
                     ),
+                    Return(
+                        value: None,
+                    ),
                 ],
                 diagnostic_filter_leaf: None,
             ),

naga/tests/out/ir/const_assert.compact.ron

@@ -43,7 +43,11 @@
             named_expressions: {
                 0: "z",
             },
-            body: [],
+            body: [
+                Return(
+                    value: None,
+                ),
+            ],
             diagnostic_filter_leaf: None,
         ),
     ],

naga/tests/out/ir/const_assert.ron

@@ -43,7 +43,11 @@
             named_expressions: {
                 0: "z",
             },
-            body: [],
+            body: [
+                Return(
+                    value: None,
+                ),
+            ],
             diagnostic_filter_leaf: None,
         ),
     ],

naga/tests/out/ir/diagnostic-filter.compact.ron

@@ -17,7 +17,11 @@
             local_variables: [],
             expressions: [],
             named_expressions: {},
-            body: [],
+            body: [
+                Return(
+                    value: None,
+                ),
+            ],
             diagnostic_filter_leaf: Some(0),
         ),
         (
@@ -27,7 +31,11 @@
             local_variables: [],
             expressions: [],
             named_expressions: {},
-            body: [],
+            body: [
+                Return(
+                    value: None,
+                ),
+            ],
             diagnostic_filter_leaf: Some(1),
         ),
     ],

naga/tests/out/ir/diagnostic-filter.ron

@@ -17,7 +17,11 @@
             local_variables: [],
             expressions: [],
             named_expressions: {},
-            body: [],
+            body: [
+                Return(
+                    value: None,
+                ),
+            ],
             diagnostic_filter_leaf: Some(0),
         ),
         (
@@ -27,7 +31,11 @@
             local_variables: [],
             expressions: [],
             named_expressions: {},
-            body: [],
+            body: [
+                Return(
+                    value: None,
+                ),
+            ],
             diagnostic_filter_leaf: Some(1),
         ),
     ],

naga/tests/out/ir/local-const.compact.ron

@@ -132,6 +132,9 @@
                     start: 4,
                     end: 5,
                 )),
+                Return(
+                    value: None,
+                ),
             ],
             diagnostic_filter_leaf: None,
         ),

naga/tests/out/ir/local-const.ron

@@ -132,6 +132,9 @@
                     start: 4,
                     end: 5,
                 )),
+                Return(
+                    value: None,
+                ),
             ],
             diagnostic_filter_leaf: None,
         ),

naga/tests/out/ir/must-use.compact.ron

@@ -146,6 +146,9 @@
                     arguments: [],
                     result: Some(0),
                 ),
+                Return(
+                    value: None,
+                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -164,7 +167,11 @@
                 local_variables: [],
                 expressions: [],
                 named_expressions: {},
-                body: [],
+                body: [
+                    Return(
+                        value: None,
+                    ),
+                ],
                 diagnostic_filter_leaf: None,
             ),
         ),

naga/tests/out/ir/must-use.ron

@@ -146,6 +146,9 @@
                     arguments: [],
                     result: Some(0),
                 ),
+                Return(
+                    value: None,
+                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -164,7 +167,11 @@
                 local_variables: [],
                 expressions: [],
                 named_expressions: {},
-                body: [],
+                body: [
+                    Return(
+                        value: None,
+                    ),
+                ],
                 diagnostic_filter_leaf: None,
             ),
         ),

naga/tests/out/ir/overrides-atomicCompareExchangeWeak.compact.ron

@@ -119,6 +119,9 @@
                         value: 3,
                         result: Some(4),
                     ),
+                    Return(
+                        value: None,
+                    ),
                 ],
                 diagnostic_filter_leaf: None,
             ),

naga/tests/out/ir/overrides-atomicCompareExchangeWeak.ron

@@ -119,6 +119,9 @@
                         value: 3,
                         result: Some(4),
                     ),
+                    Return(
+                        value: None,
+                    ),
                 ],
                 diagnostic_filter_leaf: None,
             ),

naga/tests/out/ir/overrides-ray-query.compact.ron

@@ -250,6 +250,9 @@
                         continuing: [],
                         break_if: None,
                     ),
+                    Return(
+                        value: None,
+                    ),
                 ],
                 diagnostic_filter_leaf: None,
             ),

naga/tests/out/ir/overrides-ray-query.ron

@@ -250,6 +250,9 @@
                         continuing: [],
                         break_if: None,
                     ),
+                    Return(
+                        value: None,
+                    ),
                 ],
                 diagnostic_filter_leaf: None,
             ),

naga/tests/out/ir/overrides.compact.ron

@@ -204,6 +204,9 @@
                         pointer: 11,
                         value: 12,
                     ),
+                    Return(
+                        value: None,
+                    ),
                 ],
                 diagnostic_filter_leaf: None,
             ),

naga/tests/out/ir/overrides.ron

@@ -204,6 +204,9 @@
                         pointer: 11,
                         value: 12,
                     ),
+                    Return(
+                        value: None,
+                    ),
                 ],
                 diagnostic_filter_leaf: None,
             ),

naga/tests/out/msl/6772-unpack-expr-accesses.msl

@@ -9,4 +9,5 @@ kernel void main_(
 ) {
     int phony = (int4(12u, 12u >> 8, 12u >> 16, 12u >> 24) << 24 >> 24)[2];
     uint phony_1 = (uint4(12u, 12u >> 8, 12u >> 16, 12u >> 24) << 24 >> 24).y;
+    return;
 }

naga/tests/out/msl/abstract-types-function-calls.msl

@@ -17,57 +17,68 @@ struct type_9 {
 void func_f(
     float a
 ) {
+    return;
 }
 
 void func_i(
     int a_1
 ) {
+    return;
 }
 
 void func_u(
     uint a_2
 ) {
+    return;
 }
 
 void func_vf(
     metal::float2 a_3
 ) {
+    return;
 }
 
 void func_vi(
     metal::int2 a_4
 ) {
+    return;
 }
 
 void func_vu(
     metal::uint2 a_5
 ) {
+    return;
 }
 
 void func_mf(
     metal::float2x2 a_6
 ) {
+    return;
 }
 
 void func_af(
     type_7 a_7
 ) {
+    return;
 }
 
 void func_ai(
     type_8 a_8
 ) {
+    return;
 }
 
 void func_au(
     type_9 a_9
 ) {
+    return;
 }
 
 void func_f_i(
     float a_10,
     int b
 ) {
+    return;
 }
 
 void main_(
@@ -88,4 +99,5 @@ void main_(
     func_au(type_9 {0u, 0u});
     func_f_i(0.0, 0);
     func_f_i(0.0, 0);
+    return;
 }

naga/tests/out/msl/abstract-types-operators.msl

@@ -86,14 +86,17 @@ void runtime_values(
     uint _e52 = u;
     uint _e53 = u;
     plus_u_u_u = _e52 + _e53;
+    return;
 }
 
 void wgpu_4445_(
 ) {
+    return;
 }
 
 void wgpu_4435_(
     threadgroup type_3& a
 ) {
     uint y = a.inner[as_type<int>(as_type<uint>(1) - as_type<uint>(1))];
+    return;
 }