diff --git a/CHANGELOG.md b/CHANGELOG.md
index 07315f2932..01dd3fa20f 100644
--- a/CHANGELOG.md
+++ b/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).
diff --git a/naga/src/front/glsl/functions.rs b/naga/src/front/glsl/functions.rs
index 7eccca7c1f..0d05c5433c 100644
--- a/naga/src/front/glsl/functions.rs
+++ b/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();
 
diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs
index f1e71e2c09..d25eb362c1 100644
--- a/naga/src/front/wgsl/lower/mod.rs
+++ b/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
diff --git a/naga/tests/out/glsl/6772-unpack-expr-accesses.main.Compute.glsl b/naga/tests/out/glsl/6772-unpack-expr-accesses.main.Compute.glsl
index e857de73fe..a91356afaf 100644
--- a/naga/tests/out/glsl/6772-unpack-expr-accesses.main.Compute.glsl
+++ b/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;
 }
 
diff --git a/naga/tests/out/glsl/access.assign_through_ptr.Compute.glsl b/naga/tests/out/glsl/access.assign_through_ptr.Compute.glsl
index d269990a66..004720bfc3 100644
--- a/naga/tests/out/glsl/access.assign_through_ptr.Compute.glsl
+++ b/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;
 }
 
diff --git a/naga/tests/out/glsl/access.assign_to_ptr_components.Compute.glsl b/naga/tests/out/glsl/access.assign_to_ptr_components.Compute.glsl
index 2835553237..b0224010aa 100644
--- a/naga/tests/out/glsl/access.assign_to_ptr_components.Compute.glsl
+++ b/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;
 }
 
diff --git a/naga/tests/out/glsl/access.foo_frag.Fragment.glsl b/naga/tests/out/glsl/access.foo_frag.Fragment.glsl
index 48333b727a..308cf195c1 100644
--- a/naga/tests/out/glsl/access.foo_frag.Fragment.glsl
+++ b/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) {
diff --git a/naga/tests/out/glsl/access.foo_vert.Vertex.glsl b/naga/tests/out/glsl/access.foo_vert.Vertex.glsl
index 870e603bc3..c710043d23 100644
--- a/naga/tests/out/glsl/access.foo_vert.Vertex.glsl
+++ b/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) {
diff --git a/naga/tests/out/glsl/array-in-ctor.cs_main.Compute.glsl b/naga/tests/out/glsl/array-in-ctor.cs_main.Compute.glsl
index bd918087b8..71d64585ad 100644
--- a/naga/tests/out/glsl/array-in-ctor.cs_main.Compute.glsl
+++ b/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;
 }
 
diff --git a/naga/tests/out/glsl/atomicOps.cs_main.Compute.glsl b/naga/tests/out/glsl/atomicOps.cs_main.Compute.glsl
index 64a6a9e931..b69c5107ce 100644
--- a/naga/tests/out/glsl/atomicOps.cs_main.Compute.glsl
+++ b/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;
 }
 
diff --git a/naga/tests/out/glsl/atomicTexture.cs_main.Compute.glsl b/naga/tests/out/glsl/atomicTexture.cs_main.Compute.glsl
index 993bdfea67..4defe425f9 100644
--- a/naga/tests/out/glsl/atomicTexture.cs_main.Compute.glsl
+++ b/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;
 }
 
diff --git a/naga/tests/out/glsl/bitcast.main.Compute.glsl b/naga/tests/out/glsl/bitcast.main.Compute.glsl
index 90d42e9e96..57e7e221b7 100644
--- a/naga/tests/out/glsl/bitcast.main.Compute.glsl
+++ b/naga/tests/out/glsl/bitcast.main.Compute.glsl
@@ -34,5 +34,6 @@ void main() {
     f3_ = intBitsToFloat(_e41);
     ivec4 _e43 = i4_;
     f4_ = intBitsToFloat(_e43);
+    return;
 }
 
diff --git a/naga/tests/out/glsl/bits.main.Compute.glsl b/naga/tests/out/glsl/bits.main.Compute.glsl
index 8558e6b402..f4b5c0f481 100644
--- a/naga/tests/out/glsl/bits.main.Compute.glsl
+++ b/naga/tests/out/glsl/bits.main.Compute.glsl
@@ -129,5 +129,6 @@ void main() {
     u3_ = bitfieldReverse(_e168);
     uvec4 _e170 = u4_;
     u4_ = bitfieldReverse(_e170);
+    return;
 }
 
diff --git a/naga/tests/out/glsl/boids.main.Compute.glsl b/naga/tests/out/glsl/boids.main.Compute.glsl
index b9db93cdb2..c42358bfef 100644
--- a/naga/tests/out/glsl/boids.main.Compute.glsl
+++ b/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;
 }
 
diff --git a/naga/tests/out/glsl/bounds-check-image-restrict.fragment_shader.Fragment.glsl b/naga/tests/out/glsl/bounds-check-image-restrict.fragment_shader.Fragment.glsl
index af6bad24bd..f582e85ec7 100644
--- a/naga/tests/out/glsl/bounds-check-image-restrict.fragment_shader.Fragment.glsl
+++ b/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() {
diff --git a/naga/tests/out/glsl/bounds-check-image-rzsw.fragment_shader.Fragment.glsl b/naga/tests/out/glsl/bounds-check-image-rzsw.fragment_shader.Fragment.glsl
index c7365c9046..b2096add79 100644
--- a/naga/tests/out/glsl/bounds-check-image-rzsw.fragment_shader.Fragment.glsl
+++ b/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() {
diff --git a/naga/tests/out/glsl/break-if.main.Compute.glsl b/naga/tests/out/glsl/break-if.main.Compute.glsl
index 1bcfe54006..ba96de4879 100644
--- a/naga/tests/out/glsl/break-if.main.Compute.glsl
+++ b/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;
 }
 
diff --git a/naga/tests/out/glsl/const-exprs.main.Compute.glsl b/naga/tests/out/glsl/const-exprs.main.Compute.glsl
index 7e6e9ad993..0b318a65e3 100644
--- a/naga/tests/out/glsl/const-exprs.main.Compute.glsl
+++ b/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;
 }
 
diff --git a/naga/tests/out/glsl/constructors.main.Compute.glsl b/naga/tests/out/glsl/constructors.main.Compute.glsl
index c28401d0b4..cff178f200 100644
--- a/naga/tests/out/glsl/constructors.main.Compute.glsl
+++ b/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;
 }
 
diff --git a/naga/tests/out/glsl/control-flow.main.Compute.glsl b/naga/tests/out/glsl/control-flow.main.Compute.glsl
index 4e9f63c405..391fca84f4 100644
--- a/naga/tests/out/glsl/control-flow.main.Compute.glsl
+++ b/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;
         }
     }
 }
diff --git a/naga/tests/out/glsl/cross.main.Compute.glsl b/naga/tests/out/glsl/cross.main.Compute.glsl
index a4950274e3..c78b6ebeed 100644
--- a/naga/tests/out/glsl/cross.main.Compute.glsl
+++ b/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;
 }
 
diff --git a/naga/tests/out/glsl/empty.main.Compute.glsl b/naga/tests/out/glsl/empty.main.Compute.glsl
index 6e113755b9..f6b9108cc6 100644
--- a/naga/tests/out/glsl/empty.main.Compute.glsl
+++ b/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;
 }
 
diff --git a/naga/tests/out/glsl/f64.main.Compute.glsl b/naga/tests/out/glsl/f64.main.Compute.glsl
index 42ed66e573..055965d153 100644
--- a/naga/tests/out/glsl/f64.main.Compute.glsl
+++ b/naga/tests/out/glsl/f64.main.Compute.glsl
@@ -14,5 +14,6 @@ double f(double x) {
 
 void main() {
     double _e1 = f(6.0LF);
+    return;
 }
 
diff --git a/naga/tests/out/glsl/functions-webgl.main.Fragment.glsl b/naga/tests/out/glsl/functions-webgl.main.Fragment.glsl
index bd7dea43bd..9dd084c32a 100644
--- a/naga/tests/out/glsl/functions-webgl.main.Fragment.glsl
+++ b/naga/tests/out/glsl/functions-webgl.main.Fragment.glsl
@@ -13,5 +13,6 @@ vec2 test_fma() {
 
 void main() {
     vec2 _e0 = test_fma();
+    return;
 }
 
diff --git a/naga/tests/out/glsl/functions.main.Compute.glsl b/naga/tests/out/glsl/functions.main.Compute.glsl
index ce5e351429..b0f23125f7 100644
--- a/naga/tests/out/glsl/functions.main.Compute.glsl
+++ b/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;
 }
 
diff --git a/naga/tests/out/glsl/globals.main.Compute.glsl b/naga/tests/out/glsl/globals.main.Compute.glsl
index a72ebabdcc..bc20d6415f 100644
--- a/naga/tests/out/glsl/globals.main.Compute.glsl
+++ b/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;
 }
 
diff --git a/naga/tests/out/glsl/image.main.Compute.glsl b/naga/tests/out/glsl/image.main.Compute.glsl
index 2a38f35f37..13a4b24d6e 100644
--- a/naga/tests/out/glsl/image.main.Compute.glsl
+++ b/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;
 }
 
diff --git a/naga/tests/out/glsl/interpolate_compat.frag_main.Fragment.glsl b/naga/tests/out/glsl/interpolate_compat.frag_main.Fragment.glsl
index 047d0b56b3..ac7fad324f 100644
--- a/naga/tests/out/glsl/interpolate_compat.frag_main.Fragment.glsl
+++ b/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;
 }
 
diff --git a/naga/tests/out/glsl/invariant.fs.Fragment.glsl b/naga/tests/out/glsl/invariant.fs.Fragment.glsl
index 724ff9bb3d..9936a28ad3 100644
--- a/naga/tests/out/glsl/invariant.fs.Fragment.glsl
+++ b/naga/tests/out/glsl/invariant.fs.Fragment.glsl
@@ -6,5 +6,6 @@ precision highp int;
 
 void main() {
     vec4 position = gl_FragCoord;
+    return;
 }
 
diff --git a/naga/tests/out/glsl/math-functions.main.Fragment.glsl b/naga/tests/out/glsl/math-functions.main.Fragment.glsl
index 6ac3a4de0c..1877fb3d91 100644
--- a/naga/tests/out/glsl/math-functions.main.Fragment.glsl
+++ b/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;
 }
 
diff --git a/naga/tests/out/glsl/multiview.main.Fragment.glsl b/naga/tests/out/glsl/multiview.main.Fragment.glsl
index 775d5609bc..466aea062f 100644
--- a/naga/tests/out/glsl/multiview.main.Fragment.glsl
+++ b/naga/tests/out/glsl/multiview.main.Fragment.glsl
@@ -7,5 +7,6 @@ precision highp int;
 
 void main() {
     int view_index = gl_ViewIndex;
+    return;
 }
 
diff --git a/naga/tests/out/glsl/multiview_webgl.main.Fragment.glsl b/naga/tests/out/glsl/multiview_webgl.main.Fragment.glsl
index 50e785a434..30515289c9 100644
--- a/naga/tests/out/glsl/multiview_webgl.main.Fragment.glsl
+++ b/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;
 }
 
diff --git a/naga/tests/out/glsl/operators.main.Compute.glsl b/naga/tests/out/glsl/operators.main.Compute.glsl
index e11c26156c..1c2824495f 100644
--- a/naga/tests/out/glsl/operators.main.Compute.glsl
+++ b/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;
 }
 
diff --git a/naga/tests/out/glsl/overrides.main.Compute.glsl b/naga/tests/out/glsl/overrides.main.Compute.glsl
index b1c631cec0..684e7cca58 100644
--- a/naga/tests/out/glsl/overrides.main.Compute.glsl
+++ b/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;
 }
 
diff --git a/naga/tests/out/glsl/phony_assignment.main.Compute.glsl b/naga/tests/out/glsl/phony_assignment.main.Compute.glsl
index 6118e9b4a8..1bf2a8f394 100644
--- a/naga/tests/out/glsl/phony_assignment.main.Compute.glsl
+++ b/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;
 }
 
diff --git a/naga/tests/out/glsl/separate-entry-points.compute.Compute.glsl b/naga/tests/out/glsl/separate-entry-points.compute.Compute.glsl
index 03d40c7a6c..869b7ca418 100644
--- a/naga/tests/out/glsl/separate-entry-points.compute.Compute.glsl
+++ b/naga/tests/out/glsl/separate-entry-points.compute.Compute.glsl
@@ -11,9 +11,11 @@ void barriers() {
     barrier();
     memoryBarrierShared();
     barrier();
+    return;
 }
 
 void main() {
     barriers();
+    return;
 }
 
diff --git a/naga/tests/out/glsl/separate-entry-points.fragment.Fragment.glsl b/naga/tests/out/glsl/separate-entry-points.fragment.Fragment.glsl
index 9ea32684cd..52a29a3fc2 100644
--- a/naga/tests/out/glsl/separate-entry-points.fragment.Fragment.glsl
+++ b/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() {
diff --git a/naga/tests/out/glsl/struct-layout.needs_padding_comp.Compute.glsl b/naga/tests/out/glsl/struct-layout.needs_padding_comp.Compute.glsl
index 8e682aed72..7fe97f2982 100644
--- a/naga/tests/out/glsl/struct-layout.needs_padding_comp.Compute.glsl
+++ b/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;
 }
 
diff --git a/naga/tests/out/glsl/struct-layout.no_padding_comp.Compute.glsl b/naga/tests/out/glsl/struct-layout.no_padding_comp.Compute.glsl
index 6e8a367c89..c7bbcf7b70 100644
--- a/naga/tests/out/glsl/struct-layout.no_padding_comp.Compute.glsl
+++ b/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;
 }
 
diff --git a/naga/tests/out/glsl/subgroup-operations.main.Compute.glsl b/naga/tests/out/glsl/subgroup-operations.main.Compute.glsl
index ef3eb5ac95..05ab403565 100644
--- a/naga/tests/out/glsl/subgroup-operations.main.Compute.glsl
+++ b/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;
 }
 
diff --git a/naga/tests/out/glsl/variations.main.Fragment.glsl b/naga/tests/out/glsl/variations.main.Fragment.glsl
index 1651a5e805..25b258987c 100644
--- a/naga/tests/out/glsl/variations.main.Fragment.glsl
+++ b/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() {
diff --git a/naga/tests/out/glsl/workgroup-uniform-load.test_workgroupUniformLoad.Compute.glsl b/naga/tests/out/glsl/workgroup-uniform-load.test_workgroupUniformLoad.Compute.glsl
index ee50e2e1b6..6315309c99 100644
--- a/naga/tests/out/glsl/workgroup-uniform-load.test_workgroupUniformLoad.Compute.glsl
+++ b/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;
     }
 }
 
diff --git a/naga/tests/out/glsl/workgroup-var-init.main.Compute.glsl b/naga/tests/out/glsl/workgroup-var-init.main.Compute.glsl
index 783afdaa52..de136c1109 100644
--- a/naga/tests/out/glsl/workgroup-var-init.main.Compute.glsl
+++ b/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;
 }
 
diff --git a/naga/tests/out/hlsl/6772-unpack-expr-accesses.hlsl b/naga/tests/out/hlsl/6772-unpack-expr-accesses.hlsl
index e50f40c8fc..c6d6c3f150 100644
--- a/naga/tests/out/hlsl/6772-unpack-expr-accesses.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/hlsl/access.hlsl b/naga/tests/out/hlsl/access.hlsl
index 9452f354fd..dcb62ae2af 100644
--- a/naga/tests/out/hlsl/access.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/hlsl/array-in-ctor.hlsl b/naga/tests/out/hlsl/array-in-ctor.hlsl
index 1079262a01..d7054a1899 100644
--- a/naga/tests/out/hlsl/array-in-ctor.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/hlsl/atomicOps-int64-min-max.hlsl b/naga/tests/out/hlsl/atomicOps-int64-min-max.hlsl
index 81653010b2..989a52b78b 100644
--- a/naga/tests/out/hlsl/atomicOps-int64-min-max.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/hlsl/atomicOps-int64.hlsl b/naga/tests/out/hlsl/atomicOps-int64.hlsl
index 34af783426..ea88f81753 100644
--- a/naga/tests/out/hlsl/atomicOps-int64.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/hlsl/atomicOps.hlsl b/naga/tests/out/hlsl/atomicOps.hlsl
index 29588c4353..640972a2fa 100644
--- a/naga/tests/out/hlsl/atomicOps.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/hlsl/atomicTexture-int64.hlsl b/naga/tests/out/hlsl/atomicTexture-int64.hlsl
index 19f087b34a..056489c790 100644
--- a/naga/tests/out/hlsl/atomicTexture-int64.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/hlsl/atomicTexture.hlsl b/naga/tests/out/hlsl/atomicTexture.hlsl
index 40f5789585..241cdab678 100644
--- a/naga/tests/out/hlsl/atomicTexture.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/hlsl/bitcast.hlsl b/naga/tests/out/hlsl/bitcast.hlsl
index 4e33e7aafa..b21408dda5 100644
--- a/naga/tests/out/hlsl/bitcast.hlsl
+++ b/naga/tests/out/hlsl/bitcast.hlsl
@@ -29,4 +29,5 @@ void main()
     f3_ = asfloat(_e41);
     int4 _e43 = i4_;
     f4_ = asfloat(_e43);
+    return;
 }
diff --git a/naga/tests/out/hlsl/bits.hlsl b/naga/tests/out/hlsl/bits.hlsl
index 75adce1636..1987e15514 100644
--- a/naga/tests/out/hlsl/bits.hlsl
+++ b/naga/tests/out/hlsl/bits.hlsl
@@ -300,4 +300,5 @@ void main()
     u3_ = reversebits(_e168);
     uint4 _e170 = u4_;
     u4_ = reversebits(_e170);
+    return;
 }
diff --git a/naga/tests/out/hlsl/boids.hlsl b/naga/tests/out/hlsl/boids.hlsl
index 0f881f949a..22e9c6cefd 100644
--- a/naga/tests/out/hlsl/boids.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/hlsl/bounds-check-dynamic-buffer.hlsl b/naga/tests/out/hlsl/bounds-check-dynamic-buffer.hlsl
index 969300ab8e..9e92c1ed6d 100644
--- a/naga/tests/out/hlsl/bounds-check-dynamic-buffer.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/hlsl/break-if.hlsl b/naga/tests/out/hlsl/break-if.hlsl
index ee8719776d..63a0185583 100644
--- a/naga/tests/out/hlsl/break-if.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/hlsl/collatz.hlsl b/naga/tests/out/hlsl/collatz.hlsl
index bfea10d3c7..b00586aa4c 100644
--- a/naga/tests/out/hlsl/collatz.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/hlsl/const-exprs.hlsl b/naga/tests/out/hlsl/const-exprs.hlsl
index 5f049f1b2f..aa2ba75ed6 100644
--- a/naga/tests/out/hlsl/const-exprs.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/hlsl/constructors.hlsl b/naga/tests/out/hlsl/constructors.hlsl
index 90d8db9a33..c354737114 100644
--- a/naga/tests/out/hlsl/constructors.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/hlsl/control-flow.hlsl b/naga/tests/out/hlsl/control-flow.hlsl
index d7333eadeb..2438858a8a 100644
--- a/naga/tests/out/hlsl/control-flow.hlsl
+++ b/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;
         }
     }
 }
diff --git a/naga/tests/out/hlsl/cross.hlsl b/naga/tests/out/hlsl/cross.hlsl
index 96696c5066..8d78e02e49 100644
--- a/naga/tests/out/hlsl/cross.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/hlsl/empty.hlsl b/naga/tests/out/hlsl/empty.hlsl
index e9f2b4acd0..e79d36b9c8 100644
--- a/naga/tests/out/hlsl/empty.hlsl
+++ b/naga/tests/out/hlsl/empty.hlsl
@@ -1,4 +1,5 @@
 [numthreads(1, 1, 1)]
 void main()
 {
+    return;
 }
diff --git a/naga/tests/out/hlsl/f64.hlsl b/naga/tests/out/hlsl/f64.hlsl
index 57a74cadc9..7b0fb9f67e 100644
--- a/naga/tests/out/hlsl/f64.hlsl
+++ b/naga/tests/out/hlsl/f64.hlsl
@@ -15,4 +15,5 @@ double f(double x)
 void main()
 {
     const double _e1 = f(6.0L);
+    return;
 }
diff --git a/naga/tests/out/hlsl/functions.hlsl b/naga/tests/out/hlsl/functions.hlsl
index b765e59754..6d7e210307 100644
--- a/naga/tests/out/hlsl/functions.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/hlsl/globals.hlsl b/naga/tests/out/hlsl/globals.hlsl
index 2f713eeb4b..d320899b89 100644
--- a/naga/tests/out/hlsl/globals.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/hlsl/image.hlsl b/naga/tests/out/hlsl/image.hlsl
index d749d1a688..30bfd86aeb 100644
--- a/naga/tests/out/hlsl/image.hlsl
+++ b/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)]
diff --git a/naga/tests/out/hlsl/int64.hlsl b/naga/tests/out/hlsl/int64.hlsl
index 9d60a8acba..26441380b1 100644
--- a/naga/tests/out/hlsl/int64.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/hlsl/interface.hlsl b/naga/tests/out/hlsl/interface.hlsl
index 30e2c45b75..6187ca0974 100644
--- a/naga/tests/out/hlsl/interface.hlsl
+++ b/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
diff --git a/naga/tests/out/hlsl/interpolate.hlsl b/naga/tests/out/hlsl/interpolate.hlsl
index c2ba2313ee..ebb6e83477 100644
--- a/naga/tests/out/hlsl/interpolate.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/hlsl/interpolate_compat.hlsl b/naga/tests/out/hlsl/interpolate_compat.hlsl
index 7caf19904f..85f1bb001c 100644
--- a/naga/tests/out/hlsl/interpolate_compat.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/hlsl/math-functions.hlsl b/naga/tests/out/hlsl/math-functions.hlsl
index a94b5062a6..b8ace15b3c 100644
--- a/naga/tests/out/hlsl/math-functions.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/hlsl/operators.hlsl b/naga/tests/out/hlsl/operators.hlsl
index 94e25661f6..a4a381f7e7 100644
--- a/naga/tests/out/hlsl/operators.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/hlsl/overrides.hlsl b/naga/tests/out/hlsl/overrides.hlsl
index 0b1f7852eb..5cb0e75687 100644
--- a/naga/tests/out/hlsl/overrides.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/hlsl/phony_assignment.hlsl b/naga/tests/out/hlsl/phony_assignment.hlsl
index f448290e0e..dbe587e5c8 100644
--- a/naga/tests/out/hlsl/phony_assignment.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/hlsl/ray-query.hlsl b/naga/tests/out/hlsl/ray-query.hlsl
index f4aa54b2e1..9a0a2da1ce 100644
--- a/naga/tests/out/hlsl/ray-query.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/hlsl/struct-layout.hlsl b/naga/tests/out/hlsl/struct-layout.hlsl
index cc0ac76b0c..59f046ded6 100644
--- a/naga/tests/out/hlsl/struct-layout.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/hlsl/subgroup-operations.hlsl b/naga/tests/out/hlsl/subgroup-operations.hlsl
index ab92f3cc36..839b1fa6b2 100644
--- a/naga/tests/out/hlsl/subgroup-operations.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/hlsl/workgroup-uniform-load.hlsl b/naga/tests/out/hlsl/workgroup-uniform-load.hlsl
index b8bfa44930..a8dba16b9a 100644
--- a/naga/tests/out/hlsl/workgroup-uniform-load.hlsl
+++ b/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;
     }
 }
diff --git a/naga/tests/out/hlsl/workgroup-var-init.hlsl b/naga/tests/out/hlsl/workgroup-var-init.hlsl
index 917f0b8845..49b4fe621a 100644
--- a/naga/tests/out/hlsl/workgroup-var-init.hlsl
+++ b/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;
 }
diff --git a/naga/tests/out/ir/access.compact.ron b/naga/tests/out/ir/access.compact.ron
index 8e05e2936c..53ba1a3c9b 100644
--- a/naga/tests/out/ir/access.compact.ron
+++ b/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,
             ),
diff --git a/naga/tests/out/ir/access.ron b/naga/tests/out/ir/access.ron
index 8e05e2936c..53ba1a3c9b 100644
--- a/naga/tests/out/ir/access.ron
+++ b/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,
             ),
diff --git a/naga/tests/out/ir/collatz.compact.ron b/naga/tests/out/ir/collatz.compact.ron
index d4612a7599..6a7aebe544 100644
--- a/naga/tests/out/ir/collatz.compact.ron
+++ b/naga/tests/out/ir/collatz.compact.ron
@@ -325,6 +325,9 @@
                         pointer: 4,
                         value: 10,
                     ),
+                    Return(
+                        value: None,
+                    ),
                 ],
                 diagnostic_filter_leaf: None,
             ),
diff --git a/naga/tests/out/ir/collatz.ron b/naga/tests/out/ir/collatz.ron
index d4612a7599..6a7aebe544 100644
--- a/naga/tests/out/ir/collatz.ron
+++ b/naga/tests/out/ir/collatz.ron
@@ -325,6 +325,9 @@
                         pointer: 4,
                         value: 10,
                     ),
+                    Return(
+                        value: None,
+                    ),
                 ],
                 diagnostic_filter_leaf: None,
             ),
diff --git a/naga/tests/out/ir/const_assert.compact.ron b/naga/tests/out/ir/const_assert.compact.ron
index cc2f6aa035..03c540b601 100644
--- a/naga/tests/out/ir/const_assert.compact.ron
+++ b/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,
         ),
     ],
diff --git a/naga/tests/out/ir/const_assert.ron b/naga/tests/out/ir/const_assert.ron
index cc2f6aa035..03c540b601 100644
--- a/naga/tests/out/ir/const_assert.ron
+++ b/naga/tests/out/ir/const_assert.ron
@@ -43,7 +43,11 @@
             named_expressions: {
                 0: "z",
             },
-            body: [],
+            body: [
+                Return(
+                    value: None,
+                ),
+            ],
             diagnostic_filter_leaf: None,
         ),
     ],
diff --git a/naga/tests/out/ir/diagnostic-filter.compact.ron b/naga/tests/out/ir/diagnostic-filter.compact.ron
index df8f1c9161..315c8b3423 100644
--- a/naga/tests/out/ir/diagnostic-filter.compact.ron
+++ b/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),
         ),
     ],
diff --git a/naga/tests/out/ir/diagnostic-filter.ron b/naga/tests/out/ir/diagnostic-filter.ron
index df8f1c9161..315c8b3423 100644
--- a/naga/tests/out/ir/diagnostic-filter.ron
+++ b/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),
         ),
     ],
diff --git a/naga/tests/out/ir/local-const.compact.ron b/naga/tests/out/ir/local-const.compact.ron
index 6c35ee90bc..fc3dc81e71 100644
--- a/naga/tests/out/ir/local-const.compact.ron
+++ b/naga/tests/out/ir/local-const.compact.ron
@@ -132,6 +132,9 @@
                     start: 4,
                     end: 5,
                 )),
+                Return(
+                    value: None,
+                ),
             ],
             diagnostic_filter_leaf: None,
         ),
diff --git a/naga/tests/out/ir/local-const.ron b/naga/tests/out/ir/local-const.ron
index 6c35ee90bc..fc3dc81e71 100644
--- a/naga/tests/out/ir/local-const.ron
+++ b/naga/tests/out/ir/local-const.ron
@@ -132,6 +132,9 @@
                     start: 4,
                     end: 5,
                 )),
+                Return(
+                    value: None,
+                ),
             ],
             diagnostic_filter_leaf: None,
         ),
diff --git a/naga/tests/out/ir/must-use.compact.ron b/naga/tests/out/ir/must-use.compact.ron
index c29feff15e..3d51cb0c95 100644
--- a/naga/tests/out/ir/must-use.compact.ron
+++ b/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,
             ),
         ),
diff --git a/naga/tests/out/ir/must-use.ron b/naga/tests/out/ir/must-use.ron
index c29feff15e..3d51cb0c95 100644
--- a/naga/tests/out/ir/must-use.ron
+++ b/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,
             ),
         ),
diff --git a/naga/tests/out/ir/overrides-atomicCompareExchangeWeak.compact.ron b/naga/tests/out/ir/overrides-atomicCompareExchangeWeak.compact.ron
index 6dd480e0b0..56be2f8ab6 100644
--- a/naga/tests/out/ir/overrides-atomicCompareExchangeWeak.compact.ron
+++ b/naga/tests/out/ir/overrides-atomicCompareExchangeWeak.compact.ron
@@ -119,6 +119,9 @@
                         value: 3,
                         result: Some(4),
                     ),
+                    Return(
+                        value: None,
+                    ),
                 ],
                 diagnostic_filter_leaf: None,
             ),
diff --git a/naga/tests/out/ir/overrides-atomicCompareExchangeWeak.ron b/naga/tests/out/ir/overrides-atomicCompareExchangeWeak.ron
index 6dd480e0b0..56be2f8ab6 100644
--- a/naga/tests/out/ir/overrides-atomicCompareExchangeWeak.ron
+++ b/naga/tests/out/ir/overrides-atomicCompareExchangeWeak.ron
@@ -119,6 +119,9 @@
                         value: 3,
                         result: Some(4),
                     ),
+                    Return(
+                        value: None,
+                    ),
                 ],
                 diagnostic_filter_leaf: None,
             ),
diff --git a/naga/tests/out/ir/overrides-ray-query.compact.ron b/naga/tests/out/ir/overrides-ray-query.compact.ron
index 5d4964ae67..10cad83538 100644
--- a/naga/tests/out/ir/overrides-ray-query.compact.ron
+++ b/naga/tests/out/ir/overrides-ray-query.compact.ron
@@ -250,6 +250,9 @@
                         continuing: [],
                         break_if: None,
                     ),
+                    Return(
+                        value: None,
+                    ),
                 ],
                 diagnostic_filter_leaf: None,
             ),
diff --git a/naga/tests/out/ir/overrides-ray-query.ron b/naga/tests/out/ir/overrides-ray-query.ron
index 5d4964ae67..10cad83538 100644
--- a/naga/tests/out/ir/overrides-ray-query.ron
+++ b/naga/tests/out/ir/overrides-ray-query.ron
@@ -250,6 +250,9 @@
                         continuing: [],
                         break_if: None,
                     ),
+                    Return(
+                        value: None,
+                    ),
                 ],
                 diagnostic_filter_leaf: None,
             ),
diff --git a/naga/tests/out/ir/overrides.compact.ron b/naga/tests/out/ir/overrides.compact.ron
index 8fd68b84f6..00c57fa434 100644
--- a/naga/tests/out/ir/overrides.compact.ron
+++ b/naga/tests/out/ir/overrides.compact.ron
@@ -204,6 +204,9 @@
                         pointer: 11,
                         value: 12,
                     ),
+                    Return(
+                        value: None,
+                    ),
                 ],
                 diagnostic_filter_leaf: None,
             ),
diff --git a/naga/tests/out/ir/overrides.ron b/naga/tests/out/ir/overrides.ron
index 8fd68b84f6..00c57fa434 100644
--- a/naga/tests/out/ir/overrides.ron
+++ b/naga/tests/out/ir/overrides.ron
@@ -204,6 +204,9 @@
                         pointer: 11,
                         value: 12,
                     ),
+                    Return(
+                        value: None,
+                    ),
                 ],
                 diagnostic_filter_leaf: None,
             ),
diff --git a/naga/tests/out/msl/6772-unpack-expr-accesses.msl b/naga/tests/out/msl/6772-unpack-expr-accesses.msl
index e00a1b4c54..0097ef6145 100644
--- a/naga/tests/out/msl/6772-unpack-expr-accesses.msl
+++ b/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;
 }
diff --git a/naga/tests/out/msl/abstract-types-function-calls.msl b/naga/tests/out/msl/abstract-types-function-calls.msl
index c3c6d812c8..bf62d87a54 100644
--- a/naga/tests/out/msl/abstract-types-function-calls.msl
+++ b/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;
 }
diff --git a/naga/tests/out/msl/abstract-types-operators.msl b/naga/tests/out/msl/abstract-types-operators.msl
index 633164432b..6bb3ba4380 100644
--- a/naga/tests/out/msl/abstract-types-operators.msl
+++ b/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;
 }
diff --git a/naga/tests/out/msl/abstract-types-var.msl b/naga/tests/out/msl/abstract-types-var.msl
index ed84db393d..1c9d6fcb25 100644
--- a/naga/tests/out/msl/abstract-types-var.msl
+++ b/naga/tests/out/msl/abstract-types-var.msl
@@ -87,6 +87,7 @@ void all_constant_arguments(
     iafpaiaf = type_7 {1.0, 2.0};
     iafpafai = type_7 {1.0, 2.0};
     iafpafaf = type_7 {1.0, 2.0};
+    return;
 }
 
 void mixed_constant_and_runtime_arguments(
@@ -194,4 +195,5 @@ void mixed_constant_and_runtime_arguments(
     xaip_iai = type_8 {_e169, 2};
     int _e172 = i;
     xaipai_i = type_8 {1, _e172};
+    return;
 }
diff --git a/naga/tests/out/msl/access.msl b/naga/tests/out/msl/access.msl
index e5a477517e..eb0808f80e 100644
--- a/naga/tests/out/msl/access.msl
+++ b/naga/tests/out/msl/access.msl
@@ -110,6 +110,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(
@@ -147,6 +148,7 @@ void test_matrix_within_array_within_struct_accesses(
     int _e100 = idx_1;
     int _e102 = idx_1;
     t_1.am.inner[0][_e100][_e102] = 40.0;
+    return;
 }
 
 float read_from_private(
@@ -166,12 +168,14 @@ void assign_through_ptr_fn(
     thread uint& p
 ) {
     p = 42u;
+    return;
 }
 
 void assign_array_through_ptr_fn(
     thread type_22& foo_2
 ) {
     foo_2 = type_22 {metal::float4(1.0), metal::float4(2.0)};
+    return;
 }
 
 uint fetch_arg_ptr_member(
@@ -185,6 +189,7 @@ void assign_to_arg_ptr_member(
     thread AssignToMember& p_2
 ) {
     p_2.x = 10u;
+    return;
 }
 
 uint fetch_arg_ptr_array_element(
@@ -198,6 +203,7 @@ void assign_to_arg_ptr_array_element(
     thread type_25& p_4
 ) {
     p_4.inner[1] = 10u;
+    return;
 }
 
 bool index_ptr(
@@ -298,6 +304,7 @@ kernel void assign_through_ptr(
     type_22 arr = type_22 {metal::float4(6.0), metal::float4(7.0)};
     assign_through_ptr_fn(val);
     assign_array_through_ptr_fn(arr);
+    return;
 }
 
 
@@ -309,4 +316,5 @@ kernel void assign_to_ptr_components(
     uint _e1 = fetch_arg_ptr_member(s1_);
     assign_to_arg_ptr_array_element(a1_);
     uint _e3 = fetch_arg_ptr_array_element(a1_);
+    return;
 }
diff --git a/naga/tests/out/msl/array-in-ctor.msl b/naga/tests/out/msl/array-in-ctor.msl
index a3bbb2057c..776a7b4bc6 100644
--- a/naga/tests/out/msl/array-in-ctor.msl
+++ b/naga/tests/out/msl/array-in-ctor.msl
@@ -15,4 +15,5 @@ kernel void cs_main(
   device Ah const& ah [[user(fake0)]]
 ) {
     Ah ah_1 = ah;
+    return;
 }
diff --git a/naga/tests/out/msl/atomicCompareExchange.msl b/naga/tests/out/msl/atomicCompareExchange.msl
index 7ec7eb272e..6655fad7e2 100644
--- a/naga/tests/out/msl/atomicCompareExchange.msl
+++ b/naga/tests/out/msl/atomicCompareExchange.msl
@@ -114,6 +114,7 @@ kernel void test_atomic_compare_exchange_i32_(
         }
         LOOP_IS_BOUNDED
     }
+    return;
 }
 
 
@@ -160,4 +161,5 @@ kernel void test_atomic_compare_exchange_u32_(
         }
         LOOP_IS_BOUNDED
     }
+    return;
 }
diff --git a/naga/tests/out/msl/atomicOps-float32.msl b/naga/tests/out/msl/atomicOps-float32.msl
index d062ff5852..40b5a7a14f 100644
--- a/naga/tests/out/msl/atomicOps-float32.msl
+++ b/naga/tests/out/msl/atomicOps-float32.msl
@@ -39,4 +39,5 @@ kernel void cs_main(
     float _e47 = metal::atomic_exchange_explicit(&storage_atomic_arr.inner[1], 1.5, metal::memory_order_relaxed);
     float _e51 = metal::atomic_exchange_explicit(&storage_struct.atomic_scalar, 1.5, metal::memory_order_relaxed);
     float _e56 = metal::atomic_exchange_explicit(&storage_struct.atomic_arr.inner[1], 1.5, metal::memory_order_relaxed);
+    return;
 }
diff --git a/naga/tests/out/msl/atomicOps-int64-min-max.msl b/naga/tests/out/msl/atomicOps-int64-min-max.msl
index 998f54d275..f69a2a49bd 100644
--- a/naga/tests/out/msl/atomicOps-int64-min-max.msl
+++ b/naga/tests/out/msl/atomicOps-int64-min-max.msl
@@ -34,4 +34,5 @@ kernel void cs_main(
     metal::atomic_min_explicit(&storage_atomic_arr.inner[1], 1uL + _e24, metal::memory_order_relaxed);
     metal::atomic_min_explicit(&storage_struct.atomic_scalar, 1uL, metal::memory_order_relaxed);
     metal::atomic_min_explicit(&storage_struct.atomic_arr.inner[1], static_cast<ulong>(id.x), metal::memory_order_relaxed);
+    return;
 }
diff --git a/naga/tests/out/msl/atomicOps.msl b/naga/tests/out/msl/atomicOps.msl
index 8d44e9d64c..4732b4a32d 100644
--- a/naga/tests/out/msl/atomicOps.msl
+++ b/naga/tests/out/msl/atomicOps.msl
@@ -122,4 +122,5 @@ kernel void cs_main(
     int _e295 = metal::atomic_exchange_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed);
     uint _e299 = metal::atomic_exchange_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed);
     int _e304 = metal::atomic_exchange_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed);
+    return;
 }
diff --git a/naga/tests/out/msl/atomicTexture-int64.msl b/naga/tests/out/msl/atomicTexture-int64.msl
index 59704d05ff..c00d8b7654 100644
--- a/naga/tests/out/msl/atomicTexture-int64.msl
+++ b/naga/tests/out/msl/atomicTexture-int64.msl
@@ -14,4 +14,5 @@ kernel void cs_main(
     image.atomic_max(metal::uint2(metal::int2(0, 0)), 1uL);
     metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
     image.atomic_min(metal::uint2(metal::int2(0, 0)), 1uL);
+    return;
 }
diff --git a/naga/tests/out/msl/atomicTexture.msl b/naga/tests/out/msl/atomicTexture.msl
index f61636f4e4..2cb17a081e 100644
--- a/naga/tests/out/msl/atomicTexture.msl
+++ b/naga/tests/out/msl/atomicTexture.msl
@@ -24,4 +24,5 @@ kernel void cs_main(
     image_s.atomic_fetch_and(metal::uint2(metal::int2(0, 0)), 1);
     image_s.atomic_fetch_or(metal::uint2(metal::int2(0, 0)), 1);
     image_s.atomic_fetch_xor(metal::uint2(metal::int2(0, 0)), 1);
+    return;
 }
diff --git a/naga/tests/out/msl/bitcast.msl b/naga/tests/out/msl/bitcast.msl
index 80be97f2f3..20f4b850e3 100644
--- a/naga/tests/out/msl/bitcast.msl
+++ b/naga/tests/out/msl/bitcast.msl
@@ -34,4 +34,5 @@ kernel void main_(
     f3_ = as_type<metal::float3>(_e41);
     metal::int4 _e43 = i4_;
     f4_ = as_type<metal::float4>(_e43);
+    return;
 }
diff --git a/naga/tests/out/msl/bits.msl b/naga/tests/out/msl/bits.msl
index 2a4d51a391..a0138862d5 100644
--- a/naga/tests/out/msl/bits.msl
+++ b/naga/tests/out/msl/bits.msl
@@ -129,4 +129,5 @@ kernel void main_(
     u3_ = metal::reverse_bits(_e168);
     metal::uint4 _e170 = u4_;
     u4_ = metal::reverse_bits(_e170);
+    return;
 }
diff --git a/naga/tests/out/msl/boids.msl b/naga/tests/out/msl/boids.msl
index be74f8fb45..a4d709bd09 100644
--- a/naga/tests/out/msl/boids.msl
+++ b/naga/tests/out/msl/boids.msl
@@ -156,4 +156,5 @@ kernel void main_(
     particlesDst.particles[index].pos = _e174;
     metal::float2 _e179 = vVel;
     particlesDst.particles[index].vel = _e179;
+    return;
 }
diff --git a/naga/tests/out/msl/bounds-check-image-restrict.msl b/naga/tests/out/msl/bounds-check-image-restrict.msl
index 5da5ca711f..138c0f6455 100644
--- a/naga/tests/out/msl/bounds-check-image-restrict.msl
+++ b/naga/tests/out/msl/bounds-check-image-restrict.msl
@@ -112,6 +112,7 @@ void test_textureStore_1d(
     metal::texture1d<float, metal::access::write> image_storage_1d
 ) {
     image_storage_1d.write(value, uint(coords_10));
+    return;
 }
 
 void test_textureStore_2d(
@@ -120,6 +121,7 @@ void test_textureStore_2d(
     metal::texture2d<float, metal::access::write> image_storage_2d
 ) {
     image_storage_2d.write(value_1, metal::uint2(coords_11));
+    return;
 }
 
 void test_textureStore_2d_array_u(
@@ -129,6 +131,7 @@ void test_textureStore_2d_array_u(
     metal::texture2d_array<float, metal::access::write> image_storage_2d_array
 ) {
     image_storage_2d_array.write(value_2, metal::uint2(coords_12), array_index);
+    return;
 }
 
 void test_textureStore_2d_array_s(
@@ -138,6 +141,7 @@ void test_textureStore_2d_array_s(
     metal::texture2d_array<float, metal::access::write> image_storage_2d_array
 ) {
     image_storage_2d_array.write(value_3, metal::uint2(coords_13), array_index_1);
+    return;
 }
 
 void test_textureStore_3d(
@@ -146,6 +150,7 @@ void test_textureStore_3d(
     metal::texture3d<float, metal::access::write> image_storage_3d
 ) {
     image_storage_3d.write(value_4, metal::uint3(coords_14));
+    return;
 }
 
 struct fragment_shaderOutput {
diff --git a/naga/tests/out/msl/bounds-check-image-rzsw.msl b/naga/tests/out/msl/bounds-check-image-rzsw.msl
index 8eeb535535..f73b8e3e32 100644
--- a/naga/tests/out/msl/bounds-check-image-rzsw.msl
+++ b/naga/tests/out/msl/bounds-check-image-rzsw.msl
@@ -111,6 +111,7 @@ void test_textureStore_1d(
     metal::texture1d<float, metal::access::write> image_storage_1d
 ) {
     image_storage_1d.write(value, uint(coords_10));
+    return;
 }
 
 void test_textureStore_2d(
@@ -119,6 +120,7 @@ void test_textureStore_2d(
     metal::texture2d<float, metal::access::write> image_storage_2d
 ) {
     image_storage_2d.write(value_1, metal::uint2(coords_11));
+    return;
 }
 
 void test_textureStore_2d_array_u(
@@ -128,6 +130,7 @@ void test_textureStore_2d_array_u(
     metal::texture2d_array<float, metal::access::write> image_storage_2d_array
 ) {
     image_storage_2d_array.write(value_2, metal::uint2(coords_12), array_index);
+    return;
 }
 
 void test_textureStore_2d_array_s(
@@ -137,6 +140,7 @@ void test_textureStore_2d_array_s(
     metal::texture2d_array<float, metal::access::write> image_storage_2d_array
 ) {
     image_storage_2d_array.write(value_3, metal::uint2(coords_13), array_index_1);
+    return;
 }
 
 void test_textureStore_3d(
@@ -145,6 +149,7 @@ void test_textureStore_3d(
     metal::texture3d<float, metal::access::write> image_storage_3d
 ) {
     image_storage_3d.write(value_4, metal::uint3(coords_14));
+    return;
 }
 
 struct fragment_shaderOutput {
diff --git a/naga/tests/out/msl/bounds-check-restrict.msl b/naga/tests/out/msl/bounds-check-restrict.msl
index 6202d878c1..a8bf371555 100644
--- a/naga/tests/out/msl/bounds-check-restrict.msl
+++ b/naga/tests/out/msl/bounds-check-restrict.msl
@@ -99,6 +99,7 @@ void set_array(
     constant _mslBufferSizes& _buffer_sizes
 ) {
     globals.a.inner[metal::min(unsigned(i_7), 9u)] = v_1;
+    return;
 }
 
 void set_dynamic_array(
@@ -108,6 +109,7 @@ void set_dynamic_array(
     constant _mslBufferSizes& _buffer_sizes
 ) {
     globals.d[metal::min(unsigned(i_8), (_buffer_sizes.size0 - 112 - 4) / 4)] = v_2;
+    return;
 }
 
 void set_vector(
@@ -117,6 +119,7 @@ void set_vector(
     constant _mslBufferSizes& _buffer_sizes
 ) {
     globals.v[metal::min(unsigned(i_9), 3u)] = v_3;
+    return;
 }
 
 void set_matrix(
@@ -126,6 +129,7 @@ void set_matrix(
     constant _mslBufferSizes& _buffer_sizes
 ) {
     globals.m[metal::min(unsigned(i_10), 2u)] = v_4;
+    return;
 }
 
 void set_index_twice(
@@ -136,6 +140,7 @@ void set_index_twice(
     constant _mslBufferSizes& _buffer_sizes
 ) {
     globals.m[metal::min(unsigned(i_11), 2u)][metal::min(unsigned(j_1), 3u)] = v_5;
+    return;
 }
 
 void set_expensive(
@@ -145,6 +150,7 @@ void set_expensive(
     constant _mslBufferSizes& _buffer_sizes
 ) {
     globals.a.inner[metal::min(unsigned(static_cast<int>(metal::sin(static_cast<float>(i_12) / 100.0) * 100.0)), 9u)] = v_6;
+    return;
 }
 
 void set_in_bounds(
@@ -155,6 +161,7 @@ void set_in_bounds(
     globals.a.inner[9] = v_7;
     globals.v.w = v_7;
     globals.m[2].w = v_7;
+    return;
 }
 
 float index_dynamic_array_constant_index(
@@ -171,4 +178,5 @@ void set_dynamic_array_constant_index(
     constant _mslBufferSizes& _buffer_sizes
 ) {
     globals.d[metal::min(unsigned(1000), (_buffer_sizes.size0 - 112 - 4) / 4)] = v_8;
+    return;
 }
diff --git a/naga/tests/out/msl/bounds-check-zero.msl b/naga/tests/out/msl/bounds-check-zero.msl
index 14cac0d4ef..8269d4bf70 100644
--- a/naga/tests/out/msl/bounds-check-zero.msl
+++ b/naga/tests/out/msl/bounds-check-zero.msl
@@ -108,6 +108,7 @@ void set_array(
     if (uint(i_7) < 10) {
         globals.a.inner[i_7] = v_1;
     }
+    return;
 }
 
 void set_dynamic_array(
@@ -119,6 +120,7 @@ void set_dynamic_array(
     if (uint(i_8) < 1 + (_buffer_sizes.size0 - 112 - 4) / 4) {
         globals.d[i_8] = v_2;
     }
+    return;
 }
 
 void set_vector(
@@ -130,6 +132,7 @@ void set_vector(
     if (uint(i_9) < 4) {
         globals.v[i_9] = v_3;
     }
+    return;
 }
 
 void set_matrix(
@@ -141,6 +144,7 @@ void set_matrix(
     if (uint(i_10) < 3) {
         globals.m[i_10] = v_4;
     }
+    return;
 }
 
 void set_index_twice(
@@ -153,6 +157,7 @@ void set_index_twice(
     if (uint(j_1) < 4 && uint(i_11) < 3) {
         globals.m[i_11][j_1] = v_5;
     }
+    return;
 }
 
 void set_expensive(
@@ -165,6 +170,7 @@ void set_expensive(
     if (uint(_e10) < 10) {
         globals.a.inner[_e10] = v_6;
     }
+    return;
 }
 
 void set_in_bounds(
@@ -175,6 +181,7 @@ void set_in_bounds(
     globals.a.inner[9] = v_7;
     globals.v.w = v_7;
     globals.m[2].w = v_7;
+    return;
 }
 
 float index_dynamic_array_constant_index(
@@ -193,4 +200,5 @@ void set_dynamic_array_constant_index(
     if (uint(1000) < 1 + (_buffer_sizes.size0 - 112 - 4) / 4) {
         globals.d[1000] = v_8;
     }
+    return;
 }
diff --git a/naga/tests/out/msl/break-if.msl b/naga/tests/out/msl/break-if.msl
index 5b9812436f..4d3397234b 100644
--- a/naga/tests/out/msl/break-if.msl
+++ b/naga/tests/out/msl/break-if.msl
@@ -18,6 +18,7 @@ void breakIfEmpty(
 #define LOOP_IS_BOUNDED { volatile bool unpredictable_break_from_loop = false; if (unpredictable_break_from_loop) break; }
         LOOP_IS_BOUNDED
     }
+    return;
 }
 
 void breakIfEmptyBody(
@@ -39,6 +40,7 @@ void breakIfEmptyBody(
         loop_init_1 = false;
         LOOP_IS_BOUNDED
     }
+    return;
 }
 
 void breakIf(
@@ -60,6 +62,7 @@ void breakIf(
         e = a_1 != _e2;
         LOOP_IS_BOUNDED
     }
+    return;
 }
 
 void breakIfSeparateVariable(
@@ -78,8 +81,10 @@ void breakIfSeparateVariable(
         counter = _e3 + 1u;
         LOOP_IS_BOUNDED
     }
+    return;
 }
 
 kernel void main_(
 ) {
+    return;
 }
diff --git a/naga/tests/out/msl/collatz.msl b/naga/tests/out/msl/collatz.msl
index 36dc36a25d..e282d13abb 100644
--- a/naga/tests/out/msl/collatz.msl
+++ b/naga/tests/out/msl/collatz.msl
@@ -54,4 +54,5 @@ kernel void main_(
     uint _e9 = v_indices.data[global_id.x];
     uint _e10 = collatz_iterations(_e9);
     v_indices.data[global_id.x] = _e10;
+    return;
 }
diff --git a/naga/tests/out/msl/const-exprs.msl b/naga/tests/out/msl/const-exprs.msl
index d82f5b277f..cb0959f72e 100644
--- a/naga/tests/out/msl/const-exprs.msl
+++ b/naga/tests/out/msl/const-exprs.msl
@@ -22,16 +22,19 @@ constant metal::bool2 compare_vec = metal::bool2(true, false);
 void swizzle_of_compose(
 ) {
     metal::int4 out = metal::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(
@@ -50,21 +53,25 @@ void non_constant_initializers(
     int _e10 = y;
     int _e11 = z;
     out_3 = metal::int4(_e8, _e9, _e10, _e11);
+    return;
 }
 
 void splat_of_constant(
 ) {
     metal::int4 out_4 = metal::int4(-4, -4, -4, -4);
+    return;
 }
 
 void compose_of_constant(
 ) {
     metal::int4 out_5 = metal::int4(-4, -4, -4, -4);
+    return;
 }
 
 void compose_of_splat(
 ) {
     metal::float4 x_1 = metal::float4(2.0, 1.0, 1.0, 1.0);
+    return;
 }
 
 uint map_texture_kind(
@@ -95,4 +102,5 @@ kernel void main_(
     splat_of_constant();
     compose_of_constant();
     compose_of_splat();
+    return;
 }
diff --git a/naga/tests/out/msl/constructors.msl b/naga/tests/out/msl/constructors.msl
index d4dc5c5292..1084d8e8aa 100644
--- a/naga/tests/out/msl/constructors.msl
+++ b/naga/tests/out/msl/constructors.msl
@@ -41,4 +41,5 @@ kernel void main_(
     type_11 cit2_ = type_11 {0, 1, 2, 3};
     metal::uint2 ic4_ = metal::uint2(0u, 0u);
     metal::float2x3 ic5_ = metal::float2x3(metal::float3(0.0, 0.0, 0.0), metal::float3(0.0, 0.0, 0.0));
+    return;
 }
diff --git a/naga/tests/out/msl/control-flow.msl b/naga/tests/out/msl/control-flow.msl
index 2bc471ae9e..1b35249f36 100644
--- a/naga/tests/out/msl/control-flow.msl
+++ b/naga/tests/out/msl/control-flow.msl
@@ -43,6 +43,7 @@ void loop_switch_continue(
 #define LOOP_IS_BOUNDED { volatile bool unpredictable_break_from_loop = false; if (unpredictable_break_from_loop) break; }
         LOOP_IS_BOUNDED
     }
+    return;
 }
 
 void loop_switch_continue_nesting(
@@ -102,6 +103,7 @@ void loop_switch_continue_nesting(
         }
         LOOP_IS_BOUNDED
     }
+    return;
 }
 
 void loop_switch_omit_continue_variable_checks(
@@ -154,6 +156,7 @@ void loop_switch_omit_continue_variable_checks(
         }
         LOOP_IS_BOUNDED
     }
+    return;
 }
 
 struct main_Input {
@@ -211,18 +214,18 @@ kernel void main_(
         }
         case 2: {
             pos = 1;
-            break;
+            return;
         }
         case 3: {
             pos = 2;
-            break;
+            return;
         }
         case 4: {
-            break;
+            return;
         }
         default: {
             pos = 3;
-            break;
+            return;
         }
     }
 }
diff --git a/naga/tests/out/msl/cross.msl b/naga/tests/out/msl/cross.msl
index 70095cd6e9..a402e73e32 100644
--- a/naga/tests/out/msl/cross.msl
+++ b/naga/tests/out/msl/cross.msl
@@ -8,4 +8,5 @@ using metal::uint;
 kernel void main_(
 ) {
     metal::float3 a = metal::cross(metal::float3(0.0, 1.0, 2.0), metal::float3(0.0, 1.0, 2.0));
+    return;
 }
diff --git a/naga/tests/out/msl/empty.msl b/naga/tests/out/msl/empty.msl
index 38cee7f85a..414cd22012 100644
--- a/naga/tests/out/msl/empty.msl
+++ b/naga/tests/out/msl/empty.msl
@@ -7,4 +7,5 @@ using metal::uint;
 
 kernel void main_(
 ) {
+    return;
 }
diff --git a/naga/tests/out/msl/functions.msl b/naga/tests/out/msl/functions.msl
index 7d799b52e3..42632f99be 100644
--- a/naga/tests/out/msl/functions.msl
+++ b/naga/tests/out/msl/functions.msl
@@ -31,4 +31,5 @@ kernel void main_(
 ) {
     metal::float2 _e0 = test_fma();
     int _e1 = test_integer_dot_product();
+    return;
 }
diff --git a/naga/tests/out/msl/globals.msl b/naga/tests/out/msl/globals.msl
index 4baa36d1f8..b976cc2018 100644
--- a/naga/tests/out/msl/globals.msl
+++ b/naga/tests/out/msl/globals.msl
@@ -36,6 +36,7 @@ constant bool Foo_1 = true;
 void test_msl_packed_vec3_as_arg(
     metal::float3 arg
 ) {
+    return;
 }
 
 void test_msl_packed_vec3_(
@@ -55,6 +56,7 @@ void test_msl_packed_vec3_(
     metal::float3 mvm1_ = metal::float3x3 {} * metal::float3(data.v3_);
     metal::float3 svm0_ = data.v3_ * 2.0;
     metal::float3 svm1_ = 2.0 * data.v3_;
+    return;
 }
 
 kernel void main_(
@@ -95,4 +97,5 @@ kernel void main_(
     alignment.v1_ = 4.0;
     wg.inner[1] = static_cast<float>(1 + (_buffer_sizes.size3 - 0 - 8) / 8);
     metal::atomic_store_explicit(&at_1, 2u, metal::memory_order_relaxed);
+    return;
 }
diff --git a/naga/tests/out/msl/image.msl b/naga/tests/out/msl/image.msl
index 11405d2db5..1a8567a0ca 100644
--- a/naga/tests/out/msl/image.msl
+++ b/naga/tests/out/msl/image.msl
@@ -32,6 +32,7 @@ kernel void main_(
     metal::uint4 value7u = image_1d_src.read(uint(static_cast<uint>(local_id.x)));
     image_dst.write((((value1_ + value2_) + value4_) + value5_) + value6_, uint(itc.x));
     image_dst.write((((value1u + value2u) + value4u) + value5u) + value6u, uint(static_cast<uint>(itc.x)));
+    return;
 }
 
 
diff --git a/naga/tests/out/msl/int64.msl b/naga/tests/out/msl/int64.msl
index d26c139c3c..369eb01c6f 100644
--- a/naga/tests/out/msl/int64.msl
+++ b/naga/tests/out/msl/int64.msl
@@ -209,4 +209,5 @@ kernel void main_(
     ulong _e3 = uint64_function(67uL, input_uniform, input_storage, input_arrays, output, output_arrays);
     long _e5 = int64_function(60L, input_uniform, input_storage, input_arrays, output, output_arrays);
     output.final_value = _e3 + as_type<ulong>(_e5);
+    return;
 }
diff --git a/naga/tests/out/msl/interface.msl b/naga/tests/out/msl/interface.msl
index 9dc33927e9..047873da9f 100644
--- a/naga/tests/out/msl/interface.msl
+++ b/naga/tests/out/msl/interface.msl
@@ -81,6 +81,7 @@ kernel void compute(
     }
     metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
     output.inner[0] = (((global_id.x + local_id.x) + local_index) + wg_id.x) + num_wgs.x;
+    return;
 }
 
 
diff --git a/naga/tests/out/msl/interpolate.msl b/naga/tests/out/msl/interpolate.msl
index fd3baea127..c19005753f 100644
--- a/naga/tests/out/msl/interpolate.msl
+++ b/naga/tests/out/msl/interpolate.msl
@@ -73,4 +73,5 @@ fragment void frag_main(
 , metal::float4 position [[position]]
 ) {
     const FragmentInput val = { position, varyings_1._flat, varyings_1.flat_first, varyings_1.flat_either, varyings_1._linear, varyings_1.linear_centroid, {}, varyings_1.linear_sample, varyings_1.linear_center, varyings_1.perspective, varyings_1.perspective_centroid, varyings_1.perspective_sample, varyings_1.perspective_center };
+    return;
 }
diff --git a/naga/tests/out/msl/interpolate_compat.msl b/naga/tests/out/msl/interpolate_compat.msl
index e861d71454..e386c07db0 100644
--- a/naga/tests/out/msl/interpolate_compat.msl
+++ b/naga/tests/out/msl/interpolate_compat.msl
@@ -70,4 +70,5 @@ fragment void frag_main(
 , metal::float4 position [[position]]
 ) {
     const FragmentInput val = { position, varyings_1._flat, varyings_1.flat_either, varyings_1._linear, {}, varyings_1.linear_centroid, {}, varyings_1.linear_sample, varyings_1.linear_center, varyings_1.perspective, varyings_1.perspective_centroid, varyings_1.perspective_sample, varyings_1.perspective_center };
+    return;
 }
diff --git a/naga/tests/out/msl/math-functions.msl b/naga/tests/out/msl/math-functions.msl
index f3dbe0a20d..97f3ea7317 100644
--- a/naga/tests/out/msl/math-functions.msl
+++ b/naga/tests/out/msl/math-functions.msl
@@ -93,4 +93,5 @@ fragment void main_(
     metal::float2 quantizeToF16_b = metal::float2(metal::half2(metal::float2(1.0, 1.0)));
     metal::float3 quantizeToF16_c = metal::float3(metal::half3(metal::float3(1.0, 1.0, 1.0)));
     metal::float4 quantizeToF16_d = metal::float4(metal::half4(metal::float4(1.0, 1.0, 1.0, 1.0)));
+    return;
 }
diff --git a/naga/tests/out/msl/operators.msl b/naga/tests/out/msl/operators.msl
index e6d747811c..9b9b916cb2 100644
--- a/naga/tests/out/msl/operators.msl
+++ b/naga/tests/out/msl/operators.msl
@@ -61,6 +61,7 @@ void logical(
     metal::bool3 bitwise_or1_ = metal::bool3(true) | metal::bool3(false);
     bool bitwise_and0_ = true & false;
     metal::bool4 bitwise_and1_ = metal::bool4(true) & metal::bool4(false);
+    return;
 }
 
 void arithmetic(
@@ -137,6 +138,7 @@ void arithmetic(
     metal::float3 mul_vector0_ = metal::float4x3 {} * metal::float4(1.0);
     metal::float4 mul_vector1_ = metal::float3(2.0) * metal::float4x3 {};
     metal::float3x3 mul = metal::float4x3 {} * metal::float3x4 {};
+    return;
 }
 
 void bit(
@@ -165,6 +167,7 @@ void bit(
     uint shr1_ = 2u >> 1u;
     metal::int2 shr2_ = metal::int2(2) >> metal::uint2(1u);
     metal::uint3 shr3_ = metal::uint3(2u) >> metal::uint3(1u);
+    return;
 }
 
 void comparison(
@@ -205,6 +208,7 @@ void comparison(
     metal::bool2 gte3_ = metal::int2(2) >= metal::int2(1);
     metal::bool3 gte4_ = metal::uint3(2u) >= metal::uint3(1u);
     metal::bool4 gte5_ = metal::float4(2.0) >= metal::float4(1.0);
+    return;
 }
 
 void assignment(
@@ -242,6 +246,7 @@ void assignment(
     vec0_[1] = as_type<int>(as_type<uint>(_e37) + as_type<uint>(1));
     int _e41 = vec0_[1];
     vec0_[1] = as_type<int>(as_type<uint>(_e41) - as_type<uint>(1));
+    return;
 }
 
 void negation_avoids_prefix_decrement(
@@ -254,6 +259,7 @@ void negation_avoids_prefix_decrement(
     int p5_ = -(-(-(-(1))));
     int p6_ = -(-(-(-(-(1)))));
     int p7_ = -(-(-(-(-(1)))));
+    return;
 }
 
 struct main_Input {
@@ -269,4 +275,5 @@ kernel void main_(
     bit();
     comparison();
     assignment();
+    return;
 }
diff --git a/naga/tests/out/msl/overrides-atomicCompareExchangeWeak.msl b/naga/tests/out/msl/overrides-atomicCompareExchangeWeak.msl
index dba6d8f6cb..d87190c595 100644
--- a/naga/tests/out/msl/overrides-atomicCompareExchangeWeak.msl
+++ b/naga/tests/out/msl/overrides-atomicCompareExchangeWeak.msl
@@ -44,4 +44,5 @@ kernel void f(
     }
     metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
     _atomic_compare_exchange_resultUint4_ _e5 = naga_atomic_compare_exchange_weak_explicit(&a, 2u, 1u);
+    return;
 }
diff --git a/naga/tests/out/msl/overrides-ray-query.msl b/naga/tests/out/msl/overrides-ray-query.msl
index 8dc7bb8371..3aa0ee0359 100644
--- a/naga/tests/out/msl/overrides-ray-query.msl
+++ b/naga/tests/out/msl/overrides-ray-query.msl
@@ -43,4 +43,5 @@ kernel void main_(
 #define LOOP_IS_BOUNDED { volatile bool unpredictable_break_from_loop = false; if (unpredictable_break_from_loop) break; }
         LOOP_IS_BOUNDED
     }
+    return;
 }
diff --git a/naga/tests/out/msl/overrides.msl b/naga/tests/out/msl/overrides.msl
index 32b6fd5f2e..02770d613d 100644
--- a/naga/tests/out/msl/overrides.msl
+++ b/naga/tests/out/msl/overrides.msl
@@ -24,4 +24,5 @@ kernel void main_(
     float _e9 = gain_x_10_;
     gain_x_100_ = _e9 * 10.0;
     store_override = gain;
+    return;
 }
diff --git a/naga/tests/out/msl/phony_assignment.msl b/naga/tests/out/msl/phony_assignment.msl
index daad571906..295bd7c0cc 100644
--- a/naga/tests/out/msl/phony_assignment.msl
+++ b/naga/tests/out/msl/phony_assignment.msl
@@ -21,4 +21,5 @@ kernel void main_(
     int _e6 = five();
     int _e7 = five();
     float phony_2 = binding;
+    return;
 }
diff --git a/naga/tests/out/msl/ray-query.msl b/naga/tests/out/msl/ray-query.msl
index 8d9ebbf7c5..b8230fb2e8 100644
--- a/naga/tests/out/msl/ray-query.msl
+++ b/naga/tests/out/msl/ray-query.msl
@@ -86,6 +86,7 @@ kernel void main_(
     output.visible = static_cast<uint>(_e7.kind == 0u);
     metal::float3 _e18 = get_torus_normal(dir_1 * _e7.t, _e7);
     output.normal = _e18;
+    return;
 }
 
 
@@ -104,4 +105,5 @@ kernel void main_candidate(
     rq.intersection = rq.intersector.intersect(metal::raytracing::ray(_e12.origin, _e12.dir, _e12.tmin, _e12.tmax), acc_struct, _e12.cull_mask);    rq.ready = true;
     RayIntersection intersection_1 = RayIntersection {_map_intersection_type(rq.intersection.type), rq.intersection.distance, rq.intersection.user_instance_id, rq.intersection.instance_id, {}, rq.intersection.geometry_id, rq.intersection.primitive_id, rq.intersection.triangle_barycentric_coord, rq.intersection.triangle_front_facing, {}, rq.intersection.object_to_world_transform, rq.intersection.world_to_object_transform};
     output.visible = static_cast<uint>(intersection_1.kind == 3u);
+    return;
 }
diff --git a/naga/tests/out/msl/struct-layout.msl b/naga/tests/out/msl/struct-layout.msl
index 644927b6b1..e1f266435b 100644
--- a/naga/tests/out/msl/struct-layout.msl
+++ b/naga/tests/out/msl/struct-layout.msl
@@ -54,6 +54,7 @@ kernel void no_padding_comp(
     x = _e2;
     NoPadding _e4 = no_padding_storage;
     x = _e4;
+    return;
 }
 
 
@@ -98,4 +99,5 @@ kernel void needs_padding_comp(
     x_1 = _e2;
     NeedsPadding _e4 = needs_padding_storage;
     x_1 = _e4;
+    return;
 }
diff --git a/naga/tests/out/msl/subgroup-operations.msl b/naga/tests/out/msl/subgroup-operations.msl
index e32d570e31..20a550a715 100644
--- a/naga/tests/out/msl/subgroup-operations.msl
+++ b/naga/tests/out/msl/subgroup-operations.msl
@@ -40,4 +40,5 @@ kernel void main_(
     uint unnamed_18 = metal::simd_shuffle_down(subgroup_invocation_id, 1u);
     uint unnamed_19 = metal::simd_shuffle_up(subgroup_invocation_id, 1u);
     uint unnamed_20 = metal::simd_shuffle_xor(subgroup_invocation_id, sizes.subgroup_size - 1u);
+    return;
 }
diff --git a/naga/tests/out/msl/workgroup-uniform-load.msl b/naga/tests/out/msl/workgroup-uniform-load.msl
index dfe129c37c..32495c198a 100644
--- a/naga/tests/out/msl/workgroup-uniform-load.msl
+++ b/naga/tests/out/msl/workgroup-uniform-load.msl
@@ -25,5 +25,8 @@ kernel void test_workgroupUniformLoad(
     metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
     if (unnamed > 10) {
         metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
+        return;
+    } else {
+        return;
     }
 }
diff --git a/naga/tests/out/msl/workgroup-var-init.msl b/naga/tests/out/msl/workgroup-var-init.msl
index 9d4e4d9d5c..991c8b014b 100644
--- a/naga/tests/out/msl/workgroup-var-init.msl
+++ b/naga/tests/out/msl/workgroup-var-init.msl
@@ -36,4 +36,5 @@ kernel void main_(
     metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
     type_1 _e3 = w_mem.arr;
     output = _e3;
+    return;
 }
diff --git a/naga/tests/out/spv/control-flow.spvasm b/naga/tests/out/spv/control-flow.spvasm
index 8fdb0cc34c..f3c3644b4f 100644
--- a/naga/tests/out/spv/control-flow.spvasm
+++ b/naga/tests/out/spv/control-flow.spvasm
@@ -261,15 +261,15 @@ OpStore %108 %16
 OpBranch %128
 %130 = OpLabel
 OpStore %108 %74
-OpBranch %128
+OpReturn
 %131 = OpLabel
 OpStore %108 %75
-OpBranch %128
+OpReturn
 %132 = OpLabel
-OpBranch %128
+OpReturn
 %133 = OpLabel
 OpStore %108 %105
-OpBranch %128
+OpReturn
 %128 = OpLabel
 OpReturn
 OpFunctionEnd
\ No newline at end of file
diff --git a/naga/tests/out/spv/workgroup-uniform-load.spvasm b/naga/tests/out/spv/workgroup-uniform-load.spvasm
index 69c8cc842c..87f212a799 100644
--- a/naga/tests/out/spv/workgroup-uniform-load.spvasm
+++ b/naga/tests/out/spv/workgroup-uniform-load.spvasm
@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.1
 ; Generator: rspirv
-; Bound: 39
+; Bound: 40
 OpCapability Shader
 %1 = OpExtInstImport "GLSL.std.450"
 OpMemoryModel Logical GLSL450
@@ -55,10 +55,12 @@ OpControlBarrier %29 %29 %30
 OpControlBarrier %29 %29 %30
 %36 = OpSGreaterThan  %24  %35 %16
 OpSelectionMerge %37 None
-OpBranchConditional %36 %38 %37
+OpBranchConditional %36 %38 %39
 %38 = OpLabel
 OpControlBarrier %29 %29 %30
-OpBranch %37
+OpReturn
+%39 = OpLabel
+OpReturn
 %37 = OpLabel
 OpReturn
 OpFunctionEnd
\ No newline at end of file
diff --git a/naga/tests/out/wgsl/210-bevy-2d-shader.frag.wgsl b/naga/tests/out/wgsl/210-bevy-2d-shader.frag.wgsl
index 18256ae7d2..a800b0f856 100644
--- a/naga/tests/out/wgsl/210-bevy-2d-shader.frag.wgsl
+++ b/naga/tests/out/wgsl/210-bevy-2d-shader.frag.wgsl
@@ -18,6 +18,7 @@ fn main_1() {
     color = _e4;
     let _e6 = color;
     o_Target = _e6;
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/210-bevy-2d-shader.vert.wgsl b/naga/tests/out/wgsl/210-bevy-2d-shader.vert.wgsl
index 4947398bf9..52355cd404 100644
--- a/naga/tests/out/wgsl/210-bevy-2d-shader.vert.wgsl
+++ b/naga/tests/out/wgsl/210-bevy-2d-shader.vert.wgsl
@@ -39,6 +39,7 @@ fn main_1() {
     let _e21 = global_1.Model;
     let _e23 = position;
     gl_Position = ((_e20 * _e21) * vec4<f32>(_e23.x, _e23.y, _e23.z, 1f));
+    return;
 }
 
 @vertex 
diff --git a/naga/tests/out/wgsl/210-bevy-shader.vert.wgsl b/naga/tests/out/wgsl/210-bevy-shader.vert.wgsl
index db49013db6..223538e4e0 100644
--- a/naga/tests/out/wgsl/210-bevy-shader.vert.wgsl
+++ b/naga/tests/out/wgsl/210-bevy-shader.vert.wgsl
@@ -40,6 +40,7 @@ fn main_1() {
     let _e42 = global.ViewProj;
     let _e43 = v_Position;
     gl_Position = (_e42 * vec4<f32>(_e43.x, _e43.y, _e43.z, 1f));
+    return;
 }
 
 @vertex 
diff --git a/naga/tests/out/wgsl/246-collatz.comp.wgsl b/naga/tests/out/wgsl/246-collatz.comp.wgsl
index 6fdc5106e8..71453bfe1c 100644
--- a/naga/tests/out/wgsl/246-collatz.comp.wgsl
+++ b/naga/tests/out/wgsl/246-collatz.comp.wgsl
@@ -48,6 +48,7 @@ fn main_1() {
     let _e10 = global.indices[_e8];
     let _e11 = collatz_iterations(_e10);
     global.indices[_e6] = _e11;
+    return;
 }
 
 @compute @workgroup_size(1, 1, 1) 
diff --git a/naga/tests/out/wgsl/277-casting.frag.wgsl b/naga/tests/out/wgsl/277-casting.frag.wgsl
index 9b0a19ab19..c9ef2be74a 100644
--- a/naga/tests/out/wgsl/277-casting.frag.wgsl
+++ b/naga/tests/out/wgsl/277-casting.frag.wgsl
@@ -1,6 +1,7 @@
 fn main_1() {
     var a: f32 = 1f;
 
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/280-matrix-cast.frag.wgsl b/naga/tests/out/wgsl/280-matrix-cast.frag.wgsl
index 6cc18567cd..6ac9d37540 100644
--- a/naga/tests/out/wgsl/280-matrix-cast.frag.wgsl
+++ b/naga/tests/out/wgsl/280-matrix-cast.frag.wgsl
@@ -1,6 +1,7 @@
 fn main_1() {
     var a: mat4x4<f32> = mat4x4<f32>(vec4<f32>(1f, 0f, 0f, 0f), vec4<f32>(0f, 1f, 0f, 0f), vec4<f32>(0f, 0f, 1f, 0f), vec4<f32>(0f, 0f, 0f, 1f));
 
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/484-preprocessor-if.frag.wgsl b/naga/tests/out/wgsl/484-preprocessor-if.frag.wgsl
index 152fea63be..6b3c4ac973 100644
--- a/naga/tests/out/wgsl/484-preprocessor-if.frag.wgsl
+++ b/naga/tests/out/wgsl/484-preprocessor-if.frag.wgsl
@@ -1,4 +1,5 @@
 fn main_1() {
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/5246-dual-iteration.frag.wgsl b/naga/tests/out/wgsl/5246-dual-iteration.frag.wgsl
index 315ce3d03a..e2672a9951 100644
--- a/naga/tests/out/wgsl/5246-dual-iteration.frag.wgsl
+++ b/naga/tests/out/wgsl/5246-dual-iteration.frag.wgsl
@@ -41,6 +41,7 @@ fn main_1() {
             x = (_e6 + 1i);
         }
     }
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/6772-unpack-expr-accesses.wgsl b/naga/tests/out/wgsl/6772-unpack-expr-accesses.wgsl
index 4c13bc50f3..69f8f1d48b 100644
--- a/naga/tests/out/wgsl/6772-unpack-expr-accesses.wgsl
+++ b/naga/tests/out/wgsl/6772-unpack-expr-accesses.wgsl
@@ -2,4 +2,5 @@
 fn main() {
     let phony = unpack4xI8(12u)[2i];
     let phony_1 = unpack4xU8(12u).y;
+    return;
 }
diff --git a/naga/tests/out/wgsl/800-out-of-bounds-panic.vert.wgsl b/naga/tests/out/wgsl/800-out-of-bounds-panic.vert.wgsl
index 32f57cb1fb..8d6bc50af0 100644
--- a/naga/tests/out/wgsl/800-out-of-bounds-panic.vert.wgsl
+++ b/naga/tests/out/wgsl/800-out-of-bounds-panic.vert.wgsl
@@ -29,6 +29,7 @@ fn main_1() {
     let _e20 = gl_Position;
     let _e22 = gl_Position;
     gl_Position.z = ((_e20.z + _e22.w) / 2f);
+    return;
 }
 
 @vertex 
diff --git a/naga/tests/out/wgsl/896-push-constant.frag.wgsl b/naga/tests/out/wgsl/896-push-constant.frag.wgsl
index 6c0c0c720e..729e35a43f 100644
--- a/naga/tests/out/wgsl/896-push-constant.frag.wgsl
+++ b/naga/tests/out/wgsl/896-push-constant.frag.wgsl
@@ -5,6 +5,7 @@ struct PushConstants {
 var<push_constant> c: PushConstants;
 
 fn main_1() {
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/900-implicit-conversions.frag.wgsl b/naga/tests/out/wgsl/900-implicit-conversions.frag.wgsl
index b774f07a6a..2c468dd929 100644
--- a/naga/tests/out/wgsl/900-implicit-conversions.frag.wgsl
+++ b/naga/tests/out/wgsl/900-implicit-conversions.frag.wgsl
@@ -2,54 +2,63 @@ fn exact(a: f32) {
     var a_1: f32;
 
     a_1 = a;
+    return;
 }
 
 fn exact_1(a_2: i32) {
     var a_3: i32;
 
     a_3 = a_2;
+    return;
 }
 
 fn implicit(a_4: f32) {
     var a_5: f32;
 
     a_5 = a_4;
+    return;
 }
 
 fn implicit_1(a_6: i32) {
     var a_7: i32;
 
     a_7 = a_6;
+    return;
 }
 
 fn implicit_dims(v: f32) {
     var v_1: f32;
 
     v_1 = v;
+    return;
 }
 
 fn implicit_dims_1(v_2: vec2<f32>) {
     var v_3: vec2<f32>;
 
     v_3 = v_2;
+    return;
 }
 
 fn implicit_dims_2(v_4: vec3<f32>) {
     var v_5: vec3<f32>;
 
     v_5 = v_4;
+    return;
 }
 
 fn implicit_dims_3(v_6: vec4<f32>) {
     var v_7: vec4<f32>;
 
     v_7 = v_6;
+    return;
 }
 
 fn main_1() {
     exact_1(1i);
     implicit(1f);
     implicit_dims_2(vec3(1f));
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/901-lhs-field-select.frag.wgsl b/naga/tests/out/wgsl/901-lhs-field-select.frag.wgsl
index 91c3fcf9be..ae45360074 100644
--- a/naga/tests/out/wgsl/901-lhs-field-select.frag.wgsl
+++ b/naga/tests/out/wgsl/901-lhs-field-select.frag.wgsl
@@ -2,6 +2,7 @@ fn main_1() {
     var a: vec4<f32> = vec4(1f);
 
     a.x = 2f;
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/931-constant-emitting.frag.wgsl b/naga/tests/out/wgsl/931-constant-emitting.frag.wgsl
index 284a570264..ac59a0c3ef 100644
--- a/naga/tests/out/wgsl/931-constant-emitting.frag.wgsl
+++ b/naga/tests/out/wgsl/931-constant-emitting.frag.wgsl
@@ -5,6 +5,7 @@ fn function() -> f32 {
 }
 
 fn main_1() {
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/932-for-loop-if.frag.wgsl b/naga/tests/out/wgsl/932-for-loop-if.frag.wgsl
index 74f4819de2..bf6a281b49 100644
--- a/naga/tests/out/wgsl/932-for-loop-if.frag.wgsl
+++ b/naga/tests/out/wgsl/932-for-loop-if.frag.wgsl
@@ -13,6 +13,7 @@ fn main_1() {
             i = (_e6 + 1i);
         }
     }
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/abstract-types-function-calls.wgsl b/naga/tests/out/wgsl/abstract-types-function-calls.wgsl
index 68a0f8c80f..419125bc6b 100644
--- a/naga/tests/out/wgsl/abstract-types-function-calls.wgsl
+++ b/naga/tests/out/wgsl/abstract-types-function-calls.wgsl
@@ -1,34 +1,45 @@
 fn func_f(a: f32) {
+    return;
 }
 
 fn func_i(a_1: i32) {
+    return;
 }
 
 fn func_u(a_2: u32) {
+    return;
 }
 
 fn func_vf(a_3: vec2<f32>) {
+    return;
 }
 
 fn func_vi(a_4: vec2<i32>) {
+    return;
 }
 
 fn func_vu(a_5: vec2<u32>) {
+    return;
 }
 
 fn func_mf(a_6: mat2x2<f32>) {
+    return;
 }
 
 fn func_af(a_7: array<f32, 2>) {
+    return;
 }
 
 fn func_ai(a_8: array<i32, 2>) {
+    return;
 }
 
 fn func_au(a_9: array<u32, 2>) {
+    return;
 }
 
 fn func_f_i(a_10: f32, b: i32) {
+    return;
 }
 
 fn main() {
@@ -48,5 +59,6 @@ fn main() {
     func_au(array<u32, 2>(0u, 0u));
     func_f_i(0f, 0i);
     func_f_i(0f, 0i);
+    return;
 }
 
diff --git a/naga/tests/out/wgsl/abstract-types-operators.wgsl b/naga/tests/out/wgsl/abstract-types-operators.wgsl
index 23e8046ec5..6b8ffceb14 100644
--- a/naga/tests/out/wgsl/abstract-types-operators.wgsl
+++ b/naga/tests/out/wgsl/abstract-types-operators.wgsl
@@ -79,12 +79,15 @@ fn runtime_values() {
     let _e52 = u;
     let _e53 = u;
     plus_u_u_u = (_e52 + _e53);
+    return;
 }
 
 fn wgpu_4445_() {
+    return;
 }
 
 fn wgpu_4435_() {
     let y = a[(1i - 1i)];
+    return;
 }
 
diff --git a/naga/tests/out/wgsl/abstract-types-var.wgsl b/naga/tests/out/wgsl/abstract-types-var.wgsl
index ad3b08c912..cd35501732 100644
--- a/naga/tests/out/wgsl/abstract-types-var.wgsl
+++ b/naga/tests/out/wgsl/abstract-types-var.wgsl
@@ -111,6 +111,7 @@ fn all_constant_arguments() {
     iafpaiaf = array<f32, 2>(1f, 2f);
     iafpafai = array<f32, 2>(1f, 2f);
     iafpafaf = array<f32, 2>(1f, 2f);
+    return;
 }
 
 fn mixed_constant_and_runtime_arguments() {
@@ -218,5 +219,6 @@ fn mixed_constant_and_runtime_arguments() {
     xaip_iai = array<i32, 2>(_e169, 2i);
     let _e172 = i;
     xaipai_i = array<i32, 2>(1i, _e172);
+    return;
 }
 
diff --git a/naga/tests/out/wgsl/access.wgsl b/naga/tests/out/wgsl/access.wgsl
index a43ecd8d2f..fd4f742522 100644
--- a/naga/tests/out/wgsl/access.wgsl
+++ b/naga/tests/out/wgsl/access.wgsl
@@ -84,6 +84,7 @@ fn test_matrix_within_struct_accesses() {
     let _e85 = idx;
     let _e87 = idx;
     t.m[_e85][_e87] = 40f;
+    return;
 }
 
 fn test_matrix_within_array_within_struct_accesses() {
@@ -120,6 +121,7 @@ fn test_matrix_within_array_within_struct_accesses() {
     let _e100 = idx_1;
     let _e102 = idx_1;
     t_1.am[0][_e100][_e102] = 40f;
+    return;
 }
 
 fn read_from_private(foo_1: ptr<function, f32>) -> f32 {
@@ -133,10 +135,12 @@ fn test_arr_as_arg(a: array<array<f32, 10>, 5>) -> f32 {
 
 fn assign_through_ptr_fn(p: ptr<function, u32>) {
     (*p) = 42u;
+    return;
 }
 
 fn assign_array_through_ptr_fn(foo_2: ptr<function, array<vec4<f32>, 2>>) {
     (*foo_2) = array<vec4<f32>, 2>(vec4(1f), vec4(2f));
+    return;
 }
 
 fn fetch_arg_ptr_member(p_1: ptr<function, AssignToMember>) -> u32 {
@@ -146,6 +150,7 @@ fn fetch_arg_ptr_member(p_1: ptr<function, AssignToMember>) -> u32 {
 
 fn assign_to_arg_ptr_member(p_2: ptr<function, AssignToMember>) {
     (*p_2).x = 10u;
+    return;
 }
 
 fn fetch_arg_ptr_array_element(p_3: ptr<function, array<u32, 4>>) -> u32 {
@@ -155,6 +160,7 @@ fn fetch_arg_ptr_array_element(p_3: ptr<function, array<u32, 4>>) -> u32 {
 
 fn assign_to_arg_ptr_array_element(p_4: ptr<function, array<u32, 4>>) {
     (*p_4)[1] = 10u;
+    return;
 }
 
 fn index_ptr(value: bool) -> bool {
@@ -237,6 +243,7 @@ fn assign_through_ptr() {
 
     assign_through_ptr_fn((&val));
     assign_array_through_ptr_fn((&arr));
+    return;
 }
 
 @compute @workgroup_size(1, 1, 1) 
@@ -248,4 +255,5 @@ fn assign_to_ptr_components() {
     let _e1 = fetch_arg_ptr_member((&s1_));
     assign_to_arg_ptr_array_element((&a1_));
     let _e3 = fetch_arg_ptr_array_element((&a1_));
+    return;
 }
diff --git a/naga/tests/out/wgsl/array-in-ctor.wgsl b/naga/tests/out/wgsl/array-in-ctor.wgsl
index 8c17731f0c..4cce8641b2 100644
--- a/naga/tests/out/wgsl/array-in-ctor.wgsl
+++ b/naga/tests/out/wgsl/array-in-ctor.wgsl
@@ -8,4 +8,5 @@ var<storage> ah: Ah;
 @compute @workgroup_size(1, 1, 1) 
 fn cs_main() {
     let ah_1 = ah;
+    return;
 }
diff --git a/naga/tests/out/wgsl/atomicCompareExchange-int64.wgsl b/naga/tests/out/wgsl/atomicCompareExchange-int64.wgsl
index ea4c82bac2..07cce9a89f 100644
--- a/naga/tests/out/wgsl/atomicCompareExchange-int64.wgsl
+++ b/naga/tests/out/wgsl/atomicCompareExchange-int64.wgsl
@@ -44,6 +44,7 @@ fn test_atomic_compare_exchange_i64_() {
             i = (_e26 + 1u);
         }
     }
+    return;
 }
 
 @compute @workgroup_size(1, 1, 1) 
@@ -85,4 +86,5 @@ fn test_atomic_compare_exchange_u64_() {
             i_1 = (_e26 + 1u);
         }
     }
+    return;
 }
diff --git a/naga/tests/out/wgsl/atomicCompareExchange.wgsl b/naga/tests/out/wgsl/atomicCompareExchange.wgsl
index 19aeb44251..17bd5144b1 100644
--- a/naga/tests/out/wgsl/atomicCompareExchange.wgsl
+++ b/naga/tests/out/wgsl/atomicCompareExchange.wgsl
@@ -44,6 +44,7 @@ fn test_atomic_compare_exchange_i32_() {
             i = (_e27 + 1u);
         }
     }
+    return;
 }
 
 @compute @workgroup_size(1, 1, 1) 
@@ -85,4 +86,5 @@ fn test_atomic_compare_exchange_u32_() {
             i_1 = (_e27 + 1u);
         }
     }
+    return;
 }
diff --git a/naga/tests/out/wgsl/atomicOps-float32.wgsl b/naga/tests/out/wgsl/atomicOps-float32.wgsl
index 0a4350daa2..92e67b3428 100644
--- a/naga/tests/out/wgsl/atomicOps-float32.wgsl
+++ b/naga/tests/out/wgsl/atomicOps-float32.wgsl
@@ -31,4 +31,5 @@ fn cs_main(@builtin(local_invocation_id) id: vec3<u32>) {
     let _e47 = atomicExchange((&storage_atomic_arr[1]), 1.5f);
     let _e51 = atomicExchange((&storage_struct.atomic_scalar), 1.5f);
     let _e56 = atomicExchange((&storage_struct.atomic_arr[1]), 1.5f);
+    return;
 }
diff --git a/naga/tests/out/wgsl/atomicOps-int64-min-max.wgsl b/naga/tests/out/wgsl/atomicOps-int64-min-max.wgsl
index 88e7751d1c..126758b0b5 100644
--- a/naga/tests/out/wgsl/atomicOps-int64-min-max.wgsl
+++ b/naga/tests/out/wgsl/atomicOps-int64-min-max.wgsl
@@ -27,4 +27,5 @@ fn cs_main(@builtin(local_invocation_id) id: vec3<u32>) {
     atomicMin((&storage_atomic_arr[1]), (1lu + _e24));
     atomicMin((&storage_struct.atomic_scalar), 1lu);
     atomicMin((&storage_struct.atomic_arr[1]), u64(id.x));
+    return;
 }
diff --git a/naga/tests/out/wgsl/atomicOps-int64.wgsl b/naga/tests/out/wgsl/atomicOps-int64.wgsl
index c79c8e6cae..364108c646 100644
--- a/naga/tests/out/wgsl/atomicOps-int64.wgsl
+++ b/naga/tests/out/wgsl/atomicOps-int64.wgsl
@@ -103,4 +103,5 @@ fn cs_main(@builtin(local_invocation_id) id: vec3<u32>) {
     let _e279 = atomicExchange((&workgroup_atomic_arr[1]), 1li);
     let _e283 = atomicExchange((&workgroup_struct.atomic_scalar), 1lu);
     let _e288 = atomicExchange((&workgroup_struct.atomic_arr[1]), 1li);
+    return;
 }
diff --git a/naga/tests/out/wgsl/atomicOps.wgsl b/naga/tests/out/wgsl/atomicOps.wgsl
index 696d6dd165..be102e6833 100644
--- a/naga/tests/out/wgsl/atomicOps.wgsl
+++ b/naga/tests/out/wgsl/atomicOps.wgsl
@@ -103,4 +103,5 @@ fn cs_main(@builtin(local_invocation_id) id: vec3<u32>) {
     let _e295 = atomicExchange((&workgroup_atomic_arr[1]), 1i);
     let _e299 = atomicExchange((&workgroup_struct.atomic_scalar), 1u);
     let _e304 = atomicExchange((&workgroup_struct.atomic_arr[1]), 1i);
+    return;
 }
diff --git a/naga/tests/out/wgsl/atomicTexture-int64.wgsl b/naga/tests/out/wgsl/atomicTexture-int64.wgsl
index 8cec19cbc7..52bbe76771 100644
--- a/naga/tests/out/wgsl/atomicTexture-int64.wgsl
+++ b/naga/tests/out/wgsl/atomicTexture-int64.wgsl
@@ -6,4 +6,5 @@ fn cs_main(@builtin(local_invocation_id) id: vec3<u32>) {
     textureAtomicMax(image, vec2<i32>(0i, 0i), 1lu);
     workgroupBarrier();
     textureAtomicMin(image, vec2<i32>(0i, 0i), 1lu);
+    return;
 }
diff --git a/naga/tests/out/wgsl/atomicTexture.wgsl b/naga/tests/out/wgsl/atomicTexture.wgsl
index 48c141fa2a..f5bd7db002 100644
--- a/naga/tests/out/wgsl/atomicTexture.wgsl
+++ b/naga/tests/out/wgsl/atomicTexture.wgsl
@@ -17,4 +17,5 @@ fn cs_main(@builtin(local_invocation_id) id: vec3<u32>) {
     textureAtomicAnd(image_s, vec2<i32>(0i, 0i), 1i);
     textureAtomicOr(image_s, vec2<i32>(0i, 0i), 1i);
     textureAtomicXor(image_s, vec2<i32>(0i, 0i), 1i);
+    return;
 }
diff --git a/naga/tests/out/wgsl/bevy-pbr.frag.wgsl b/naga/tests/out/wgsl/bevy-pbr.frag.wgsl
index c627967cc5..78d3fbe606 100644
--- a/naga/tests/out/wgsl/bevy-pbr.frag.wgsl
+++ b/naga/tests/out/wgsl/bevy-pbr.frag.wgsl
@@ -849,6 +849,7 @@ fn main_1() {
     output_color.z = _e301.z;
     let _e308 = output_color;
     o_Target = _e308;
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/bevy-pbr.vert.wgsl b/naga/tests/out/wgsl/bevy-pbr.vert.wgsl
index e8c85052c4..1524b9e469 100644
--- a/naga/tests/out/wgsl/bevy-pbr.vert.wgsl
+++ b/naga/tests/out/wgsl/bevy-pbr.vert.wgsl
@@ -49,6 +49,7 @@ fn main_1() {
     let _e56 = global.ViewProj;
     let _e57 = world_position;
     gl_Position = (_e56 * _e57);
+    return;
 }
 
 @vertex 
diff --git a/naga/tests/out/wgsl/bitcast.wgsl b/naga/tests/out/wgsl/bitcast.wgsl
index 1c5e31502a..1f36ce5e66 100644
--- a/naga/tests/out/wgsl/bitcast.wgsl
+++ b/naga/tests/out/wgsl/bitcast.wgsl
@@ -28,4 +28,5 @@ fn main() {
     f3_ = bitcast<vec3<f32>>(_e41);
     let _e43 = i4_;
     f4_ = bitcast<vec4<f32>>(_e43);
+    return;
 }
diff --git a/naga/tests/out/wgsl/bits.wgsl b/naga/tests/out/wgsl/bits.wgsl
index 23371c247e..0d23b1e782 100644
--- a/naga/tests/out/wgsl/bits.wgsl
+++ b/naga/tests/out/wgsl/bits.wgsl
@@ -123,4 +123,5 @@ fn main() {
     u3_ = reverseBits(_e168);
     let _e170 = u4_;
     u4_ = reverseBits(_e170);
+    return;
 }
diff --git a/naga/tests/out/wgsl/bits_glsl.frag.wgsl b/naga/tests/out/wgsl/bits_glsl.frag.wgsl
index 833398422e..3513d1ba25 100644
--- a/naga/tests/out/wgsl/bits_glsl.frag.wgsl
+++ b/naga/tests/out/wgsl/bits_glsl.frag.wgsl
@@ -102,6 +102,7 @@ fn main_1() {
     i3_ = vec3<i32>(firstLeadingBit(_e190));
     let _e193 = u4_;
     i4_ = vec4<i32>(firstLeadingBit(_e193));
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/boids.wgsl b/naga/tests/out/wgsl/boids.wgsl
index cb64da41fe..ee28ca6786 100644
--- a/naga/tests/out/wgsl/boids.wgsl
+++ b/naga/tests/out/wgsl/boids.wgsl
@@ -144,4 +144,5 @@ fn main(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {
     particlesDst.particles[index].pos = _e174;
     let _e179 = vVel;
     particlesDst.particles[index].vel = _e179;
+    return;
 }
diff --git a/naga/tests/out/wgsl/bool-select.frag.wgsl b/naga/tests/out/wgsl/bool-select.frag.wgsl
index 525c22ba6b..4045a3de45 100644
--- a/naga/tests/out/wgsl/bool-select.frag.wgsl
+++ b/naga/tests/out/wgsl/bool-select.frag.wgsl
@@ -34,6 +34,7 @@ fn main_1() {
     o_color.z = _e7.z;
     let _e17 = TevPerCompGT(3f, 5f);
     o_color.w = _e17;
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/break-if.wgsl b/naga/tests/out/wgsl/break-if.wgsl
index 5356f3d866..c3d45a50ac 100644
--- a/naga/tests/out/wgsl/break-if.wgsl
+++ b/naga/tests/out/wgsl/break-if.wgsl
@@ -4,6 +4,7 @@ fn breakIfEmpty() {
             break if true;
         }
     }
+    return;
 }
 
 fn breakIfEmptyBody(a: bool) {
@@ -19,6 +20,7 @@ fn breakIfEmptyBody(a: bool) {
             break if (a == _e5);
         }
     }
+    return;
 }
 
 fn breakIf(a_1: bool) {
@@ -34,6 +36,7 @@ fn breakIf(a_1: bool) {
             break if (a_1 == _e5);
         }
     }
+    return;
 }
 
 fn breakIfSeparateVariable() {
@@ -47,8 +50,10 @@ fn breakIfSeparateVariable() {
             break if (_e5 == 5u);
         }
     }
+    return;
 }
 
 @compute @workgroup_size(1, 1, 1) 
 fn main() {
+    return;
 }
diff --git a/naga/tests/out/wgsl/buffer.frag.wgsl b/naga/tests/out/wgsl/buffer.frag.wgsl
index 9251b978af..349d062e74 100644
--- a/naga/tests/out/wgsl/buffer.frag.wgsl
+++ b/naga/tests/out/wgsl/buffer.frag.wgsl
@@ -20,6 +20,7 @@ fn main_1() {
     testBuffer.data[1i] = 2u;
     let _e19 = testBufferReadOnly.data[0];
     b = _e19;
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/clamp-splat.vert.wgsl b/naga/tests/out/wgsl/clamp-splat.vert.wgsl
index da6f754576..909ff07209 100644
--- a/naga/tests/out/wgsl/clamp-splat.vert.wgsl
+++ b/naga/tests/out/wgsl/clamp-splat.vert.wgsl
@@ -9,6 +9,7 @@ fn main_1() {
     let _e2 = a_pos_1;
     let _e7 = clamp(_e2, vec2(0f), vec2(1f));
     gl_Position = vec4<f32>(_e7.x, _e7.y, 0f, 1f);
+    return;
 }
 
 @vertex 
diff --git a/naga/tests/out/wgsl/collatz.wgsl b/naga/tests/out/wgsl/collatz.wgsl
index c72395367f..477317594f 100644
--- a/naga/tests/out/wgsl/collatz.wgsl
+++ b/naga/tests/out/wgsl/collatz.wgsl
@@ -38,4 +38,5 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
     let _e9 = v_indices.data[global_id.x];
     let _e10 = collatz_iterations(_e9);
     v_indices.data[global_id.x] = _e10;
+    return;
 }
diff --git a/naga/tests/out/wgsl/const-exprs.wgsl b/naga/tests/out/wgsl/const-exprs.wgsl
index f8a53da0a4..411e835b5d 100644
--- a/naga/tests/out/wgsl/const-exprs.wgsl
+++ b/naga/tests/out/wgsl/const-exprs.wgsl
@@ -16,16 +16,19 @@ const compare_vec: vec2<bool> = vec2<bool>(true, false);
 fn swizzle_of_compose() {
     var out: vec4<i32> = vec4<i32>(4i, 3i, 2i, 1i);
 
+    return;
 }
 
 fn index_of_compose() {
     var out_1: i32 = 2i;
 
+    return;
 }
 
 fn compose_three_deep() {
     var out_2: i32 = 6i;
 
+    return;
 }
 
 fn non_constant_initializers() {
@@ -44,21 +47,25 @@ fn non_constant_initializers() {
     let _e10 = y;
     let _e11 = z;
     out_3 = vec4<i32>(_e8, _e9, _e10, _e11);
+    return;
 }
 
 fn splat_of_constant() {
     var out_4: vec4<i32> = vec4<i32>(-4i, -4i, -4i, -4i);
 
+    return;
 }
 
 fn compose_of_constant() {
     var out_5: vec4<i32> = vec4<i32>(-4i, -4i, -4i, -4i);
 
+    return;
 }
 
 fn compose_of_splat() {
     var x_1: vec4<f32> = vec4<f32>(2f, 1f, 1f, 1f);
 
+    return;
 }
 
 fn map_texture_kind(texture_kind: i32) -> u32 {
@@ -87,4 +94,5 @@ fn main() {
     splat_of_constant();
     compose_of_constant();
     compose_of_splat();
+    return;
 }
diff --git a/naga/tests/out/wgsl/const-global-swizzle.frag.wgsl b/naga/tests/out/wgsl/const-global-swizzle.frag.wgsl
index 5b5bf6f348..dee49c2463 100644
--- a/naga/tests/out/wgsl/const-global-swizzle.frag.wgsl
+++ b/naga/tests/out/wgsl/const-global-swizzle.frag.wgsl
@@ -14,6 +14,7 @@ fn main_1() {
     col = (_e3.xy * blank);
     let _e7 = col;
     o_Target = vec4<f32>(_e7.x, _e7.y, 0f, 1f);
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/const_assert.wgsl b/naga/tests/out/wgsl/const_assert.wgsl
index 84aceb1e80..fe1c1c02f4 100644
--- a/naga/tests/out/wgsl/const_assert.wgsl
+++ b/naga/tests/out/wgsl/const_assert.wgsl
@@ -2,5 +2,6 @@ const x: i32 = 1i;
 const y: i32 = 2i;
 
 fn foo() {
+    return;
 }
 
diff --git a/naga/tests/out/wgsl/constant-array-size.frag.wgsl b/naga/tests/out/wgsl/constant-array-size.frag.wgsl
index 56ebf9ceb7..46697cb7c2 100644
--- a/naga/tests/out/wgsl/constant-array-size.frag.wgsl
+++ b/naga/tests/out/wgsl/constant-array-size.frag.wgsl
@@ -32,6 +32,7 @@ fn function() -> vec4<f32> {
 }
 
 fn main_1() {
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/constructors.wgsl b/naga/tests/out/wgsl/constructors.wgsl
index 6d9d7e2f5d..622903d7eb 100644
--- a/naga/tests/out/wgsl/constructors.wgsl
+++ b/naga/tests/out/wgsl/constructors.wgsl
@@ -28,4 +28,5 @@ fn main() {
     const cit2_ = array<i32, 4>(0i, 1i, 2i, 3i);
     const ic4_ = vec2<u32>(0u, 0u);
     const ic5_ = mat2x3<f32>(vec3<f32>(0f, 0f, 0f), vec3<f32>(0f, 0f, 0f));
+    return;
 }
diff --git a/naga/tests/out/wgsl/control-flow.wgsl b/naga/tests/out/wgsl/control-flow.wgsl
index 69ef57740d..ad071af58a 100644
--- a/naga/tests/out/wgsl/control-flow.wgsl
+++ b/naga/tests/out/wgsl/control-flow.wgsl
@@ -27,6 +27,7 @@ fn loop_switch_continue(x: i32) {
             }
         }
     }
+    return;
 }
 
 fn loop_switch_continue_nesting(x_1: i32, y: i32, z: i32) {
@@ -73,6 +74,7 @@ fn loop_switch_continue_nesting(x_1: i32, y: i32, z: i32) {
             }
         }
     }
+    return;
 }
 
 fn loop_switch_omit_continue_variable_checks(x_2: i32, y_1: i32, z_1: i32, w: i32) {
@@ -111,6 +113,7 @@ fn loop_switch_omit_continue_variable_checks(x_2: i32, y_1: i32, z_1: i32, w: i3
             }
         }
     }
+    return;
 }
 
 @compute @workgroup_size(1, 1, 1) 
@@ -157,14 +160,18 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
         }
         case 2: {
             pos = 1i;
+            return;
         }
         case 3: {
             pos = 2i;
+            return;
         }
         case 4: {
+            return;
         }
         default: {
             pos = 3i;
+            return;
         }
     }
 }
diff --git a/naga/tests/out/wgsl/cross.wgsl b/naga/tests/out/wgsl/cross.wgsl
index 2e213aa9c4..c9ed52f5af 100644
--- a/naga/tests/out/wgsl/cross.wgsl
+++ b/naga/tests/out/wgsl/cross.wgsl
@@ -1,4 +1,5 @@
 @compute @workgroup_size(1, 1, 1) 
 fn main() {
     let a = cross(vec3<f32>(0f, 1f, 2f), vec3<f32>(0f, 1f, 2f));
+    return;
 }
diff --git a/naga/tests/out/wgsl/declarations.frag.wgsl b/naga/tests/out/wgsl/declarations.frag.wgsl
index d8f47b0b00..c0c0236ad6 100644
--- a/naga/tests/out/wgsl/declarations.frag.wgsl
+++ b/naga/tests/out/wgsl/declarations.frag.wgsl
@@ -51,6 +51,7 @@ fn main_1() {
     let _e57 = array_toomanyd[0][0][0][0][0][0][0];
     b = _e57;
     out_array[0i] = vec4(2f);
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/double-math-functions.frag.wgsl b/naga/tests/out/wgsl/double-math-functions.frag.wgsl
index 75eed3eec3..01c798586a 100644
--- a/naga/tests/out/wgsl/double-math-functions.frag.wgsl
+++ b/naga/tests/out/wgsl/double-math-functions.frag.wgsl
@@ -90,6 +90,7 @@ fn main_1() {
     smoothStepScalar = f64(smoothstep(0f, 1f, 0.5f));
     smoothStepVector = smoothstep(vec4(0.0lf), vec4(1.0lf), vec4(0.5lf));
     smoothStepMixed = smoothstep(vec4(0.0lf), vec4(1.0lf), vec4(0.5lf));
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/empty.wgsl b/naga/tests/out/wgsl/empty.wgsl
index 46b5369bef..9d1c6ac9d9 100644
--- a/naga/tests/out/wgsl/empty.wgsl
+++ b/naga/tests/out/wgsl/empty.wgsl
@@ -1,3 +1,4 @@
 @compute @workgroup_size(1, 1, 1) 
 fn main() {
+    return;
 }
diff --git a/naga/tests/out/wgsl/expressions.frag.wgsl b/naga/tests/out/wgsl/expressions.frag.wgsl
index f8632f103e..aec2b41c8b 100644
--- a/naga/tests/out/wgsl/expressions.frag.wgsl
+++ b/naga/tests/out/wgsl/expressions.frag.wgsl
@@ -36,6 +36,7 @@ fn testBinOpVecFloat(a: vec4<f32>, b: f32) {
     v = (_e12 + vec4(2f));
     let _e16 = a_1;
     v = (_e16 - vec4(2f));
+    return;
 }
 
 fn testBinOpFloatVec(a_2: vec4<f32>, b_2: f32) {
@@ -57,6 +58,7 @@ fn testBinOpFloatVec(a_2: vec4<f32>, b_2: f32) {
     let _e16 = a_3;
     let _e17 = b_3;
     v_1 = (_e16 - vec4(_e17));
+    return;
 }
 
 fn testBinOpIVecInt(a_4: vec4<i32>, b_4: i32) {
@@ -93,6 +95,7 @@ fn testBinOpIVecInt(a_4: vec4<i32>, b_4: i32) {
     let _e37 = a_5;
     let _e38 = b_5;
     v_2 = (_e37 << vec4(u32(_e38)));
+    return;
 }
 
 fn testBinOpIntIVec(a_6: i32, b_6: vec4<i32>) {
@@ -120,6 +123,7 @@ fn testBinOpIntIVec(a_6: i32, b_6: vec4<i32>) {
     let _e24 = a_7;
     let _e25 = b_7;
     v_3 = (vec4(_e24) ^ _e25);
+    return;
 }
 
 fn testBinOpUVecUint(a_8: vec4<u32>, b_8: u32) {
@@ -156,6 +160,7 @@ fn testBinOpUVecUint(a_8: vec4<u32>, b_8: u32) {
     let _e36 = a_9;
     let _e37 = b_9;
     v_4 = (_e36 << vec4(_e37));
+    return;
 }
 
 fn testBinOpUintUVec(a_10: u32, b_10: vec4<u32>) {
@@ -183,6 +188,7 @@ fn testBinOpUintUVec(a_10: u32, b_10: vec4<u32>) {
     let _e24 = a_11;
     let _e25 = b_11;
     v_5 = (vec4(_e24) ^ _e25);
+    return;
 }
 
 fn testBinOpMatMat(a_12: mat3x3<f32>, b_12: mat3x3<f32>) {
@@ -211,6 +217,7 @@ fn testBinOpMatMat(a_12: mat3x3<f32>, b_12: mat3x3<f32>) {
     let _e43 = a_13;
     let _e44 = b_13;
     c = (any((_e43[2] != _e44[2])) || (any((_e43[1] != _e44[1])) || any((_e43[0] != _e44[0]))));
+    return;
 }
 
 fn testBinOpMatFloat(a_14: f32, b_14: mat3x3<f32>) {
@@ -250,6 +257,7 @@ fn testBinOpMatFloat(a_14: f32, b_14: mat3x3<f32>) {
     let _e62 = a_15;
     let _e63 = vec3(_e62);
     v_7 = mat3x3<f32>((_e61[0] - _e63), (_e61[1] - _e63), (_e61[2] - _e63));
+    return;
 }
 
 fn testUnaryOpMat(a_16: mat3x3<f32>) {
@@ -268,24 +276,29 @@ fn testUnaryOpMat(a_16: mat3x3<f32>) {
     const _e12 = vec3(1f);
     a_17 = (_e10 - mat3x3<f32>(_e12, _e12, _e12));
     v_8 = _e10;
+    return;
 }
 
 fn testStructConstructor() {
     var tree: BST = BST(1i);
 
+    return;
 }
 
 fn testNonScalarToScalarConstructor() {
     var f: f32 = 1f;
 
+    return;
 }
 
 fn testArrayConstructor() {
     var tree_1: array<f32, 1> = array<f32, 1>(0f);
 
+    return;
 }
 
 fn testFreestandingConstructor() {
+    return;
 }
 
 fn testNonImplicitCastVectorCast() {
@@ -294,9 +307,11 @@ fn testNonImplicitCastVectorCast() {
 
     let _e3 = a_18;
     b_16 = vec4(i32(_e3));
+    return;
 }
 
 fn privatePointer(a_19: ptr<function, f32>) {
+    return;
 }
 
 fn ternary(a_20: bool) {
@@ -349,6 +364,7 @@ fn ternary(a_20: bool) {
     }
     let _e31 = local_4;
     nested = _e31;
+    return;
 }
 
 fn testMatrixMultiplication(a_22: mat4x3<f32>, b_18: mat4x4<f32>) {
@@ -361,12 +377,14 @@ fn testMatrixMultiplication(a_22: mat4x3<f32>, b_18: mat4x4<f32>) {
     let _e5 = a_23;
     let _e6 = b_19;
     c_2 = (_e5 * _e6);
+    return;
 }
 
 fn testLength() {
     var len: i32;
 
     len = i32(arrayLength((&global_1.a)));
+    return;
 }
 
 fn testConstantLength(a_24: array<f32, 4>) {
@@ -374,6 +392,7 @@ fn testConstantLength(a_24: array<f32, 4>) {
     var len_1: i32 = 4i;
 
     a_25 = a_24;
+    return;
 }
 
 fn indexConstantNonConstantIndex(i: i32) {
@@ -385,6 +404,7 @@ fn indexConstantNonConstantIndex(i: i32) {
     let _e6 = i_1;
     let _e11 = local_5.array_[_e6];
     a_26 = _e11;
+    return;
 }
 
 fn testSwizzleWrites(a_27: vec3<f32>) {
@@ -403,6 +423,7 @@ fn testSwizzleWrites(a_27: vec3<f32>) {
     let _e28 = (_e24.zy + vec2(1f));
     a_28.z = _e28.x;
     a_28.y = _e28.y;
+    return;
 }
 
 fn main_1() {
@@ -418,6 +439,7 @@ fn main_1() {
     o_color.y = 1f;
     o_color.z = 1f;
     o_color.w = 1f;
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/f64.wgsl b/naga/tests/out/wgsl/f64.wgsl
index c461908323..65699237fe 100644
--- a/naga/tests/out/wgsl/f64.wgsl
+++ b/naga/tests/out/wgsl/f64.wgsl
@@ -13,4 +13,5 @@ fn f(x: f64) -> f64 {
 @compute @workgroup_size(1, 1, 1) 
 fn main() {
     let _e1 = f(6.0lf);
+    return;
 }
diff --git a/naga/tests/out/wgsl/fma.frag.wgsl b/naga/tests/out/wgsl/fma.frag.wgsl
index 69be9f8e1b..57d11c1d24 100644
--- a/naga/tests/out/wgsl/fma.frag.wgsl
+++ b/naga/tests/out/wgsl/fma.frag.wgsl
@@ -28,6 +28,7 @@ fn Fma(d: ptr<function, Mat4x3_>, m: Mat4x3_, s: f32) {
     let _e24 = m_1;
     let _e26 = s_1;
     (*d).mz = (_e22.mz + (_e24.mz * _e26));
+    return;
 }
 
 fn main_1() {
@@ -36,6 +37,7 @@ fn main_1() {
     o_color.y = 1f;
     o_color.z = 1f;
     o_color.w = 1f;
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/functions.wgsl b/naga/tests/out/wgsl/functions.wgsl
index 0887cb513d..db7b81b146 100644
--- a/naga/tests/out/wgsl/functions.wgsl
+++ b/naga/tests/out/wgsl/functions.wgsl
@@ -20,4 +20,5 @@ fn test_integer_dot_product() -> i32 {
 fn main() {
     let _e0 = test_fma();
     let _e1 = test_integer_dot_product();
+    return;
 }
diff --git a/naga/tests/out/wgsl/functions_call.frag.wgsl b/naga/tests/out/wgsl/functions_call.frag.wgsl
index c68f5fe0b1..62aaeb2033 100644
--- a/naga/tests/out/wgsl/functions_call.frag.wgsl
+++ b/naga/tests/out/wgsl/functions_call.frag.wgsl
@@ -1,4 +1,5 @@
 fn swizzleCallee(a: ptr<function, vec2<f32>>) {
+    return;
 }
 
 fn swizzleCaller(a_1: vec3<f32>) {
@@ -14,9 +15,11 @@ fn swizzleCaller(a_1: vec3<f32>) {
     a_2.x = _e11;
     let _e12 = local.y;
     a_2.z = _e12;
+    return;
 }
 
 fn outImplicitCastCallee(a_3: ptr<function, u32>) {
+    return;
 }
 
 fn outImplicitCastCaller(a_4: f32) {
@@ -27,9 +30,11 @@ fn outImplicitCastCaller(a_4: f32) {
     outImplicitCastCallee((&local_1));
     let _e5 = local_1;
     a_5 = f32(_e5);
+    return;
 }
 
 fn swizzleImplicitCastCallee(a_6: ptr<function, vec2<u32>>) {
+    return;
 }
 
 fn swizzleImplicitCastCaller(a_7: vec3<f32>) {
@@ -44,9 +49,11 @@ fn swizzleImplicitCastCaller(a_7: vec3<f32>) {
     a_8.x = f32(_e11);
     let _e13 = local_2.y;
     a_8.z = f32(_e13);
+    return;
 }
 
 fn main_1() {
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/global-constant-array.frag.wgsl b/naga/tests/out/wgsl/global-constant-array.frag.wgsl
index bdb509dc35..0370f9b16a 100644
--- a/naga/tests/out/wgsl/global-constant-array.frag.wgsl
+++ b/naga/tests/out/wgsl/global-constant-array.frag.wgsl
@@ -6,6 +6,7 @@ fn main_1() {
     var local: array<f32, 2> = array_;
 
     let _e2 = i;
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/globals.wgsl b/naga/tests/out/wgsl/globals.wgsl
index e85185c85f..d9e2bdb2cd 100644
--- a/naga/tests/out/wgsl/globals.wgsl
+++ b/naga/tests/out/wgsl/globals.wgsl
@@ -23,6 +23,7 @@ var<uniform> global_nested_arrays_of_matrices_2x4_: array<array<mat2x4<f32>, 2>,
 var<uniform> global_nested_arrays_of_matrices_4x2_: array<array<mat4x2<f32>, 2>, 2>;
 
 fn test_msl_packed_vec3_as_arg(arg: vec3<f32>) {
+    return;
 }
 
 fn test_msl_packed_vec3_() {
@@ -41,6 +42,7 @@ fn test_msl_packed_vec3_() {
     let mvm1_ = (mat3x3<f32>() * data.v3_);
     let svm0_ = (data.v3_ * 2f);
     let svm1_ = (2f * data.v3_);
+    return;
 }
 
 @compute @workgroup_size(1, 1, 1) 
@@ -66,4 +68,5 @@ fn main() {
     alignment.v1_ = 4f;
     wg[1] = f32(arrayLength((&dummy)));
     atomicStore((&at_1), 2u);
+    return;
 }
diff --git a/naga/tests/out/wgsl/image.wgsl b/naga/tests/out/wgsl/image.wgsl
index 21647f7fef..0c10eda857 100644
--- a/naga/tests/out/wgsl/image.wgsl
+++ b/naga/tests/out/wgsl/image.wgsl
@@ -61,6 +61,7 @@ fn main(@builtin(local_invocation_id) local_id: vec3<u32>) {
     let value7u = textureLoad(image_1d_src, u32(local_id.x), i32(local_id.z));
     textureStore(image_dst, itc.x, ((((value1_ + value2_) + value4_) + value5_) + value6_));
     textureStore(image_dst, u32(itc.x), ((((value1u + value2u) + value4u) + value5u) + value6u));
+    return;
 }
 
 @compute @workgroup_size(16, 1, 1) 
diff --git a/naga/tests/out/wgsl/images.frag.wgsl b/naga/tests/out/wgsl/images.frag.wgsl
index c91bb338e1..6b57b52aeb 100644
--- a/naga/tests/out/wgsl/images.frag.wgsl
+++ b/naga/tests/out/wgsl/images.frag.wgsl
@@ -28,6 +28,7 @@ fn testImg1D(coord: i32) {
     let _e17 = coord_1;
     let _e18 = textureLoad(img1D, _e17);
     c = _e18;
+    return;
 }
 
 fn testImg1DArray(coord_2: vec2<i32>) {
@@ -44,6 +45,7 @@ fn testImg1DArray(coord_2: vec2<i32>) {
     c_1 = _e19;
     let _e21 = coord_3;
     textureStore(img1DArray, _e21.x, _e21.y, vec4(2f));
+    return;
 }
 
 fn testImg2D(coord_4: vec2<i32>) {
@@ -59,6 +61,7 @@ fn testImg2D(coord_4: vec2<i32>) {
     c_2 = _e15;
     let _e17 = coord_5;
     textureStore(img2D, _e17, vec4(2f));
+    return;
 }
 
 fn testImg2DArray(coord_6: vec3<i32>) {
@@ -75,6 +78,7 @@ fn testImg2DArray(coord_6: vec3<i32>) {
     c_3 = _e21;
     let _e23 = coord_7;
     textureStore(img2DArray, _e23.xy, _e23.z, vec4(2f));
+    return;
 }
 
 fn testImg3D(coord_8: vec3<i32>) {
@@ -90,6 +94,7 @@ fn testImg3D(coord_8: vec3<i32>) {
     c_4 = _e15;
     let _e17 = coord_9;
     textureStore(img3D, _e17, vec4(2f));
+    return;
 }
 
 fn testImgReadOnly(coord_10: vec2<i32>) {
@@ -103,6 +108,7 @@ fn testImgReadOnly(coord_10: vec2<i32>) {
     let _e14 = coord_11;
     let _e15 = textureLoad(imgReadOnly, _e14);
     c_5 = _e15;
+    return;
 }
 
 fn testImgWriteOnly(coord_12: vec2<i32>) {
@@ -114,6 +120,7 @@ fn testImgWriteOnly(coord_12: vec2<i32>) {
     size_6 = vec2<f32>(vec2<i32>(_e10));
     let _e14 = coord_13;
     textureStore(imgWriteOnly, _e14, vec4(2f));
+    return;
 }
 
 fn testImgWriteReadOnly(coord_14: vec2<i32>) {
@@ -123,9 +130,11 @@ fn testImgWriteReadOnly(coord_14: vec2<i32>) {
     coord_15 = coord_14;
     let _e10 = textureDimensions(imgWriteReadOnly);
     size_7 = vec2<f32>(vec2<i32>(_e10));
+    return;
 }
 
 fn main_1() {
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/int64.wgsl b/naga/tests/out/wgsl/int64.wgsl
index 900702812b..19d383eba5 100644
--- a/naga/tests/out/wgsl/int64.wgsl
+++ b/naga/tests/out/wgsl/int64.wgsl
@@ -186,4 +186,5 @@ fn main() {
     let _e3 = uint64_function(67lu);
     let _e5 = int64_function(60li);
     output.final_value = (_e3 + bitcast<u64>(_e5));
+    return;
 }
diff --git a/naga/tests/out/wgsl/interface.wgsl b/naga/tests/out/wgsl/interface.wgsl
index d46197662c..1f46278ab8 100644
--- a/naga/tests/out/wgsl/interface.wgsl
+++ b/naga/tests/out/wgsl/interface.wgsl
@@ -35,6 +35,7 @@ fn fragment(in: VertexOutput, @builtin(front_facing) front_facing: bool, @builti
 @compute @workgroup_size(1, 1, 1) 
 fn compute(@builtin(global_invocation_id) global_id: vec3<u32>, @builtin(local_invocation_id) local_id: vec3<u32>, @builtin(local_invocation_index) local_index: u32, @builtin(workgroup_id) wg_id: vec3<u32>, @builtin(num_workgroups) num_wgs: vec3<u32>) {
     output[0] = ((((global_id.x + local_id.x) + local_index) + wg_id.x) + num_wgs.x);
+    return;
 }
 
 @vertex 
diff --git a/naga/tests/out/wgsl/interpolate.wgsl b/naga/tests/out/wgsl/interpolate.wgsl
index c9d8a6f5c2..adb90568e4 100644
--- a/naga/tests/out/wgsl/interpolate.wgsl
+++ b/naga/tests/out/wgsl/interpolate.wgsl
@@ -35,4 +35,5 @@ fn vert_main() -> FragmentInput {
 
 @fragment 
 fn frag_main(val: FragmentInput) {
+    return;
 }
diff --git a/naga/tests/out/wgsl/interpolate_compat.wgsl b/naga/tests/out/wgsl/interpolate_compat.wgsl
index 706f9ce7fb..aeaac63547 100644
--- a/naga/tests/out/wgsl/interpolate_compat.wgsl
+++ b/naga/tests/out/wgsl/interpolate_compat.wgsl
@@ -33,4 +33,5 @@ fn vert_main() -> FragmentInput {
 
 @fragment 
 fn frag_main(val: FragmentInput) {
+    return;
 }
diff --git a/naga/tests/out/wgsl/inverse-polyfill.frag.wgsl b/naga/tests/out/wgsl/inverse-polyfill.frag.wgsl
index 325b7e7d40..1efea1d9f6 100644
--- a/naga/tests/out/wgsl/inverse-polyfill.frag.wgsl
+++ b/naga/tests/out/wgsl/inverse-polyfill.frag.wgsl
@@ -25,6 +25,7 @@ fn main_1() {
     m3_inverse = _naga_inverse_3x3_f32(_e66);
     let _e69 = m2_;
     m2_inverse = _naga_inverse_2x2_f32(_e69);
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/lexical-scopes.wgsl b/naga/tests/out/wgsl/lexical-scopes.wgsl
index fd9fa6aa92..f14fc1a8c2 100644
--- a/naga/tests/out/wgsl/lexical-scopes.wgsl
+++ b/naga/tests/out/wgsl/lexical-scopes.wgsl
@@ -1,18 +1,23 @@
 fn blockLexicalScope(a: bool) {
     {
         {
+            return;
         }
     }
 }
 
 fn ifLexicalScope(a_1: bool) {
     if a_1 {
+        return;
+    } else {
+        return;
     }
 }
 
 fn loopLexicalScope(a_2: bool) {
     loop {
     }
+    return;
 }
 
 fn forLexicalScope(a_3: f32) {
@@ -31,6 +36,7 @@ fn forLexicalScope(a_3: f32) {
             a_4 = (_e8 + 1i);
         }
     }
+    return;
 }
 
 fn whileLexicalScope(a_5: i32) {
@@ -42,6 +48,7 @@ fn whileLexicalScope(a_5: i32) {
         {
         }
     }
+    return;
 }
 
 fn switchLexicalScope(a_6: i32) {
@@ -54,5 +61,6 @@ fn switchLexicalScope(a_6: i32) {
         }
     }
     let test = (a_6 == 2i);
+    return;
 }
 
diff --git a/naga/tests/out/wgsl/local-const.wgsl b/naga/tests/out/wgsl/local-const.wgsl
index 587f5a8e54..06e52a9277 100644
--- a/naga/tests/out/wgsl/local-const.wgsl
+++ b/naga/tests/out/wgsl/local-const.wgsl
@@ -7,5 +7,6 @@ const gf: f32 = 2f;
 
 fn const_in_fn() {
     const e = vec3<i32>(4i, 4i, 4i);
+    return;
 }
 
diff --git a/naga/tests/out/wgsl/local-var-init-in-loop.comp.wgsl b/naga/tests/out/wgsl/local-var-init-in-loop.comp.wgsl
index 26f56576ae..a4bb55437a 100644
--- a/naga/tests/out/wgsl/local-var-init-in-loop.comp.wgsl
+++ b/naga/tests/out/wgsl/local-var-init-in-loop.comp.wgsl
@@ -19,6 +19,7 @@ fn main_1() {
             i = (_e10 + 1i);
         }
     }
+    return;
 }
 
 @compute @workgroup_size(1, 1, 1) 
diff --git a/naga/tests/out/wgsl/long-form-matrix.frag.wgsl b/naga/tests/out/wgsl/long-form-matrix.frag.wgsl
index c69439159a..68ccff91aa 100644
--- a/naga/tests/out/wgsl/long-form-matrix.frag.wgsl
+++ b/naga/tests/out/wgsl/long-form-matrix.frag.wgsl
@@ -8,6 +8,7 @@ fn main_1() {
     var d: mat3x3<f32> = mat3x3<f32>(vec3<f32>(2f, 2f, 1f), vec3<f32>(1f, 1f, 1f), vec3<f32>(1f, 1f, 1f));
     var e: mat4x4<f32> = mat4x4<f32>(vec4<f32>(2f, 2f, 1f, 1f), vec4<f32>(1f, 1f, 2f, 2f), vec4<f32>(1f, 1f, 1f, 1f), vec4<f32>(1f, 1f, 1f, 1f));
 
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/math-functions.frag.wgsl b/naga/tests/out/wgsl/math-functions.frag.wgsl
index 608a37a2bc..10dc84c9a5 100644
--- a/naga/tests/out/wgsl/math-functions.frag.wgsl
+++ b/naga/tests/out/wgsl/math-functions.frag.wgsl
@@ -149,6 +149,7 @@ fn main_1() {
     smoothStepScalar = smoothstep(0f, 1f, 0.5f);
     smoothStepVector = smoothstep(vec4(0f), vec4(1f), vec4(0.5f));
     smoothStepMixed = smoothstep(vec4(0f), vec4(1f), vec4(0.5f));
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/math-functions.wgsl b/naga/tests/out/wgsl/math-functions.wgsl
index f48a5dd025..b4d876a7f6 100644
--- a/naga/tests/out/wgsl/math-functions.wgsl
+++ b/naga/tests/out/wgsl/math-functions.wgsl
@@ -36,4 +36,5 @@ fn main() {
     let quantizeToF16_b = quantizeToF16(vec2<f32>(1f, 1f));
     let quantizeToF16_c = quantizeToF16(vec3<f32>(1f, 1f, 1f));
     let quantizeToF16_d = quantizeToF16(vec4<f32>(1f, 1f, 1f, 1f));
+    return;
 }
diff --git a/naga/tests/out/wgsl/module-scope.wgsl b/naga/tests/out/wgsl/module-scope.wgsl
index d9a41d3146..c9b0496a64 100644
--- a/naga/tests/out/wgsl/module-scope.wgsl
+++ b/naga/tests/out/wgsl/module-scope.wgsl
@@ -10,6 +10,7 @@ var Texture: texture_2d<f32>;
 var Sampler: sampler;
 
 fn statement() {
+    return;
 }
 
 fn returns() -> S {
@@ -20,5 +21,6 @@ fn call() {
     statement();
     let _e0 = returns();
     let s = textureSample(Texture, Sampler, vec2(1f));
+    return;
 }
 
diff --git a/naga/tests/out/wgsl/multiview.wgsl b/naga/tests/out/wgsl/multiview.wgsl
index 69dd53775c..51192d2f7a 100644
--- a/naga/tests/out/wgsl/multiview.wgsl
+++ b/naga/tests/out/wgsl/multiview.wgsl
@@ -1,3 +1,4 @@
 @fragment 
 fn main(@builtin(view_index) view_index: i32) {
+    return;
 }
diff --git a/naga/tests/out/wgsl/operators.wgsl b/naga/tests/out/wgsl/operators.wgsl
index 711b53c7a4..4a6f1c8ab4 100644
--- a/naga/tests/out/wgsl/operators.wgsl
+++ b/naga/tests/out/wgsl/operators.wgsl
@@ -48,6 +48,7 @@ fn logical() {
     let bitwise_or1_ = (vec3(true) | vec3(false));
     let bitwise_and0_ = (true & false);
     let bitwise_and1_ = (vec4(true) & vec4(false));
+    return;
 }
 
 fn arithmetic() {
@@ -123,6 +124,7 @@ fn arithmetic() {
     let mul_vector0_ = (mat4x3<f32>() * vec4(1f));
     let mul_vector1_ = (vec3(2f) * mat4x3<f32>());
     let mul = (mat4x3<f32>() * mat3x4<f32>());
+    return;
 }
 
 fn bit() {
@@ -150,6 +152,7 @@ fn bit() {
     let shr1_ = (2u >> 1u);
     let shr2_ = (vec2(2i) >> vec2(1u));
     let shr3_ = (vec3(2u) >> vec3(1u));
+    return;
 }
 
 fn comparison() {
@@ -189,6 +192,7 @@ fn comparison() {
     let gte3_ = (vec2(2i) >= vec2(1i));
     let gte4_ = (vec3(2u) >= vec3(1u));
     let gte5_ = (vec4(2f) >= vec4(1f));
+    return;
 }
 
 fn assignment() {
@@ -226,6 +230,7 @@ fn assignment() {
     vec0_[1i] = (_e37 + 1i);
     let _e41 = vec0_[1i];
     vec0_[1i] = (_e41 - 1i);
+    return;
 }
 
 fn negation_avoids_prefix_decrement() {
@@ -237,6 +242,7 @@ fn negation_avoids_prefix_decrement() {
     const p5_ = -(-(-(-(1i))));
     const p6_ = -(-(-(-(-(1i)))));
     const p7_ = -(-(-(-(-(1i)))));
+    return;
 }
 
 @compute @workgroup_size(1, 1, 1) 
@@ -249,4 +255,5 @@ fn main(@builtin(workgroup_id) id: vec3<u32>) {
     bit();
     comparison();
     assignment();
+    return;
 }
diff --git a/naga/tests/out/wgsl/phony_assignment.wgsl b/naga/tests/out/wgsl/phony_assignment.wgsl
index 60987a15cb..8e28b74314 100644
--- a/naga/tests/out/wgsl/phony_assignment.wgsl
+++ b/naga/tests/out/wgsl/phony_assignment.wgsl
@@ -12,4 +12,5 @@ fn main(@builtin(global_invocation_id) id: vec3<u32>) {
     let _e6 = five();
     let _e7 = five();
     let phony_2 = binding;
+    return;
 }
diff --git a/naga/tests/out/wgsl/pointers.wgsl b/naga/tests/out/wgsl/pointers.wgsl
index 42b0927f98..310f53cf08 100644
--- a/naga/tests/out/wgsl/pointers.wgsl
+++ b/naga/tests/out/wgsl/pointers.wgsl
@@ -10,16 +10,19 @@ fn f() {
 
     let px = (&v.x);
     (*px) = 10i;
+    return;
 }
 
 fn index_unsized(i: i32, v_1: u32) {
     let val = dynamic_array.arr[i];
     dynamic_array.arr[i] = (val + v_1);
+    return;
 }
 
 fn index_dynamic_array(i_1: i32, v_2: u32) {
     let p = (&dynamic_array.arr);
     let val_1 = (*p)[i_1];
     (*p)[i_1] = (val_1 + v_2);
+    return;
 }
 
diff --git a/naga/tests/out/wgsl/prepostfix.frag.wgsl b/naga/tests/out/wgsl/prepostfix.frag.wgsl
index 6ce6114da8..15916303bb 100644
--- a/naga/tests/out/wgsl/prepostfix.frag.wgsl
+++ b/naga/tests/out/wgsl/prepostfix.frag.wgsl
@@ -29,6 +29,7 @@ fn main_1() {
     let _e41 = (_e37 - mat4x3<f32>(_e39, _e39, _e39, _e39));
     mat = _e41;
     mat_target = _e41;
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/quad_glsl.frag.wgsl b/naga/tests/out/wgsl/quad_glsl.frag.wgsl
index 23d82a6b93..9de0a19c4b 100644
--- a/naga/tests/out/wgsl/quad_glsl.frag.wgsl
+++ b/naga/tests/out/wgsl/quad_glsl.frag.wgsl
@@ -7,6 +7,7 @@ var<private> o_color: vec4<f32>;
 
 fn main_1() {
     o_color = vec4<f32>(1f, 1f, 1f, 1f);
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/quad_glsl.vert.wgsl b/naga/tests/out/wgsl/quad_glsl.vert.wgsl
index ca7c705f9d..0a3d7cecac 100644
--- a/naga/tests/out/wgsl/quad_glsl.vert.wgsl
+++ b/naga/tests/out/wgsl/quad_glsl.vert.wgsl
@@ -16,6 +16,7 @@ fn main_1() {
     let _e6 = a_pos_1;
     let _e7 = (c_scale * _e6);
     gl_Position = vec4<f32>(_e7.x, _e7.y, 0f, 1f);
+    return;
 }
 
 @vertex 
diff --git a/naga/tests/out/wgsl/sampler-functions.frag.wgsl b/naga/tests/out/wgsl/sampler-functions.frag.wgsl
index 332134287d..0f79bd9693 100644
--- a/naga/tests/out/wgsl/sampler-functions.frag.wgsl
+++ b/naga/tests/out/wgsl/sampler-functions.frag.wgsl
@@ -27,6 +27,7 @@ fn CalcShadowPCF(T_P_t_TextureDepth_1: texture_depth_2d, S_P_t_TextureDepth_1: s
 }
 
 fn main_1() {
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/samplers.frag.wgsl b/naga/tests/out/wgsl/samplers.frag.wgsl
index 1dcece4ae2..170b0491b3 100644
--- a/naga/tests/out/wgsl/samplers.frag.wgsl
+++ b/naga/tests/out/wgsl/samplers.frag.wgsl
@@ -118,6 +118,7 @@ fn testTex1D(coord: f32) {
     let _e179 = coord_1;
     let _e183 = textureLoad(tex1D, i32(_e179), 3i);
     c = _e183;
+    return;
 }
 
 fn testTex1DArray(coord_2: vec2<f32>) {
@@ -158,6 +159,7 @@ fn testTex1DArray(coord_2: vec2<f32>) {
     let _e76 = vec2<i32>(_e75);
     let _e81 = textureLoad(tex1DArray, _e76.x, _e76.y, 3i);
     c_1 = _e81;
+    return;
 }
 
 fn testTex2D(coord_4: vec2<f32>) {
@@ -277,6 +279,7 @@ fn testTex2D(coord_4: vec2<f32>) {
     let _e307 = coord_5;
     let _e312 = textureLoad(itex2D, vec2<i32>(_e307), 3i);
     c_2 = vec4<f32>(_e312);
+    return;
 }
 
 fn testTex2DShadow(coord_6: vec2<f32>) {
@@ -344,6 +347,7 @@ fn testTex2DShadow(coord_6: vec2<f32>) {
     let _e182 = (_e176.xyz / vec3(_e176.w));
     let _e185 = textureSampleCompare(tex2DShadow, sampShadow, _e182.xy, _e182.z, vec2(5i));
     d = _e185;
+    return;
 }
 
 fn testTex2DArray(coord_8: vec3<f32>) {
@@ -390,6 +394,7 @@ fn testTex2DArray(coord_8: vec3<f32>) {
     let _e99 = vec3<i32>(_e98);
     let _e105 = textureLoad(tex2DArray, _e99.xy, _e99.z, 3i);
     c_3 = _e105;
+    return;
 }
 
 fn testTex2DArrayShadow(coord_10: vec3<f32>) {
@@ -420,6 +425,7 @@ fn testTex2DArrayShadow(coord_10: vec3<f32>) {
     let _e80 = vec4<f32>(_e75.x, _e75.y, _e75.z, 1f);
     let _e87 = textureSampleCompare(tex2DArrayShadow, sampShadow, _e80.xy, i32(_e80.z), _e80.w, vec2(5i));
     d_1 = _e87;
+    return;
 }
 
 fn testTexCube(coord_12: vec3<f32>) {
@@ -445,6 +451,7 @@ fn testTexCube(coord_12: vec3<f32>) {
     let _e39 = coord_13;
     let _e41 = textureSampleLevel(texCube, samp, _e39, 3f);
     c_4 = _e41;
+    return;
 }
 
 fn testTexCubeShadow(coord_14: vec3<f32>) {
@@ -466,6 +473,7 @@ fn testTexCubeShadow(coord_14: vec3<f32>) {
     let _e42 = vec4<f32>(_e37.x, _e37.y, _e37.z, 1f);
     let _e49 = textureSampleCompareLevel(texCubeShadow, sampShadow, _e42.xyz, _e42.w);
     d_2 = _e49;
+    return;
 }
 
 fn testTexCubeArray(coord_16: vec4<f32>) {
@@ -492,6 +500,7 @@ fn testTexCubeArray(coord_16: vec4<f32>) {
     let _e52 = coord_17;
     let _e57 = textureSampleLevel(texCubeArray, samp, _e52.xyz, i32(_e52.w), 3f);
     c_5 = _e57;
+    return;
 }
 
 fn testTexCubeArrayShadow(coord_18: vec4<f32>) {
@@ -509,6 +518,7 @@ fn testTexCubeArrayShadow(coord_18: vec4<f32>) {
     let _e32 = coord_19;
     let _e37 = textureSampleCompare(texCubeArrayShadow, sampShadow, _e32.xyz, i32(_e32.w), 1f);
     d_3 = _e37;
+    return;
 }
 
 fn testTex3D(coord_20: vec3<f32>) {
@@ -584,6 +594,7 @@ fn testTex3D(coord_20: vec3<f32>) {
     let _e176 = coord_21;
     let _e181 = textureLoad(tex3D, vec3<i32>(_e176), 3i);
     c_6 = _e181;
+    return;
 }
 
 fn testTex2DMS(coord_22: vec2<f32>) {
@@ -597,6 +608,7 @@ fn testTex2DMS(coord_22: vec2<f32>) {
     let _e24 = coord_23;
     let _e27 = textureLoad(tex2DMS, vec2<i32>(_e24), 3i);
     c_7 = _e27;
+    return;
 }
 
 fn testTex2DMSArray(coord_24: vec3<f32>) {
@@ -612,9 +624,11 @@ fn testTex2DMSArray(coord_24: vec3<f32>) {
     let _e29 = vec3<i32>(_e28);
     let _e33 = textureLoad(tex2DMSArray, _e29.xy, _e29.z, 3i);
     c_8 = _e33;
+    return;
 }
 
 fn main_1() {
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/statements.frag.wgsl b/naga/tests/out/wgsl/statements.frag.wgsl
index 8a25c1e1c2..bc36eb2075 100644
--- a/naga/tests/out/wgsl/statements.frag.wgsl
+++ b/naga/tests/out/wgsl/statements.frag.wgsl
@@ -54,6 +54,7 @@ fn switchNoLastBreak(a_6: i32) {
 }
 
 fn main_1() {
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/struct-layout.wgsl b/naga/tests/out/wgsl/struct-layout.wgsl
index e03348c4d3..b1569bde3e 100644
--- a/naga/tests/out/wgsl/struct-layout.wgsl
+++ b/naga/tests/out/wgsl/struct-layout.wgsl
@@ -36,6 +36,7 @@ fn no_padding_comp() {
     x = _e2;
     let _e4 = no_padding_storage;
     x = _e4;
+    return;
 }
 
 @fragment 
@@ -56,4 +57,5 @@ fn needs_padding_comp() {
     x_1 = _e2;
     let _e4 = needs_padding_storage;
     x_1 = _e4;
+    return;
 }
diff --git a/naga/tests/out/wgsl/subgroup-operations.wgsl b/naga/tests/out/wgsl/subgroup-operations.wgsl
index b1fa9c2511..25f713b357 100644
--- a/naga/tests/out/wgsl/subgroup-operations.wgsl
+++ b/naga/tests/out/wgsl/subgroup-operations.wgsl
@@ -27,4 +27,5 @@ fn main(sizes: Structure, @builtin(subgroup_id) subgroup_id: u32, @builtin(subgr
     let _e35 = subgroupShuffleDown(subgroup_invocation_id, 1u);
     let _e37 = subgroupShuffleUp(subgroup_invocation_id, 1u);
     let _e41 = subgroupShuffleXor(subgroup_invocation_id, (sizes.subgroup_size - 1u));
+    return;
 }
diff --git a/naga/tests/out/wgsl/type-alias.wgsl b/naga/tests/out/wgsl/type-alias.wgsl
index 13bfcba82c..ffb8dcbb08 100644
--- a/naga/tests/out/wgsl/type-alias.wgsl
+++ b/naga/tests/out/wgsl/type-alias.wgsl
@@ -6,5 +6,6 @@ fn main() {
     const e = vec3<i32>(d);
     const f = mat2x2<f32>(vec2<f32>(1f, 2f), vec2<f32>(3f, 4f));
     const g = mat3x3<f32>(a, a, a);
+    return;
 }
 
diff --git a/naga/tests/out/wgsl/vector-functions.frag.wgsl b/naga/tests/out/wgsl/vector-functions.frag.wgsl
index bb5e771411..0f28a40c87 100644
--- a/naga/tests/out/wgsl/vector-functions.frag.wgsl
+++ b/naga/tests/out/wgsl/vector-functions.frag.wgsl
@@ -28,6 +28,7 @@ fn ftest(a: vec4<f32>, b: vec4<f32>) {
     let _e24 = a_1;
     let _e25 = b_1;
     h = (_e24 != _e25);
+    return;
 }
 
 fn dtest(a_2: vec4<f64>, b_2: vec4<f64>) {
@@ -60,6 +61,7 @@ fn dtest(a_2: vec4<f64>, b_2: vec4<f64>) {
     let _e24 = a_3;
     let _e25 = b_3;
     h_1 = (_e24 != _e25);
+    return;
 }
 
 fn itest(a_4: vec4<i32>, b_4: vec4<i32>) {
@@ -92,6 +94,7 @@ fn itest(a_4: vec4<i32>, b_4: vec4<i32>) {
     let _e24 = a_5;
     let _e25 = b_5;
     h_2 = (_e24 != _e25);
+    return;
 }
 
 fn utest(a_6: vec4<u32>, b_6: vec4<u32>) {
@@ -124,6 +127,7 @@ fn utest(a_6: vec4<u32>, b_6: vec4<u32>) {
     let _e24 = a_7;
     let _e25 = b_7;
     h_3 = (_e24 != _e25);
+    return;
 }
 
 fn btest(a_8: vec4<bool>, b_8: vec4<bool>) {
@@ -149,9 +153,11 @@ fn btest(a_8: vec4<bool>, b_8: vec4<bool>) {
     f_4 = all(_e15);
     let _e18 = a_9;
     g_4 = !(_e18);
+    return;
 }
 
 fn main_1() {
+    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/workgroup-uniform-load.wgsl b/naga/tests/out/wgsl/workgroup-uniform-load.wgsl
index 10568d0e75..be85b033b9 100644
--- a/naga/tests/out/wgsl/workgroup-uniform-load.wgsl
+++ b/naga/tests/out/wgsl/workgroup-uniform-load.wgsl
@@ -8,5 +8,8 @@ fn test_workgroupUniformLoad(@builtin(workgroup_id) workgroup_id: vec3<u32>) {
     let _e4 = workgroupUniformLoad(x);
     if (_e4 > 10i) {
         workgroupBarrier();
+        return;
+    } else {
+        return;
     }
 }
diff --git a/naga/tests/out/wgsl/workgroup-var-init.wgsl b/naga/tests/out/wgsl/workgroup-var-init.wgsl
index 0b41e4bb15..fdad0477d6 100644
--- a/naga/tests/out/wgsl/workgroup-var-init.wgsl
+++ b/naga/tests/out/wgsl/workgroup-var-init.wgsl
@@ -12,4 +12,5 @@ var<storage, read_write> output: array<u32, 512>;
 fn main() {
     let _e3 = w_mem.arr;
     output = _e3;
+    return;
 }