Repository: gfx-rs/wgpu (language: rust)

Issue:
[wgsl-in] Function with missing return statement is sometimes accepted without an error
This function is accepted, even though it doesn't return anything. Seems like it should be an error.

```
fn f() -> i32 {
    let x = 1 + 1;
}
```

Strangely, it's only accepted in certain situations. If the `let` is changed to `var` or the expression is simplified to just `1` then it produces a validation error.


Discussion / hints:
I can't reproduce this on latest master but I think it was happening due to `Emit` being considered a valid return statement which should not be the case.

https://github.com/gfx-rs/wgpu/blob/4bc1bf000c27c296a587f449f2f8788b9d654657/naga/src/proc/terminator.rs#L31
The testcase was "fixed" by a730236b68f0e73f3bafbba767134751e9ea830f, but not really. Prior to this commit the function contained an `Emit` statement for the addition expression. As Teo points out, this means `ensure_block_returns` does nothing. That commit made it so that the 1 + 1 is const-evaluated away, leaving the function with no expressions to emit. This means `ensure_block_returns` adds a `return None`, which fails validation as that does not match the function return type of int.

We can still, however, reproduce the bug with this testcase:
```wgsl
fn f(x: i32) -> i32 {
  let y = x + 10;
}
```

This validates successfully, but clearly it should not.

Teo's suggestion of removing `Emit` from the list of valid return statements. instead moving it to the match arm below which makes us append a `return None`, indeed fixes the issue. It does mean we now add a useless `return;` statement to the end of some void functions. Bu we already do this anyway for the cases terminator.rs does handle (eg blocks ending in something other than an Emit)

We can easily avoid this by not calling `ensure_block_returns` for void functions. Either way the snapshots are going to have a lot of changes. I think I have a slight preference for omitting the useless return statements, but purely on aesthetic terms. So if there's even a moderately convincing argument for keeping them that's fine by me.



Fix patch:
diff --git a/naga/src/front/glsl/functions.rs b/naga/src/front/glsl/functions.rs
index 0d05c5433c..7eccca7c1f 100644
--- a/naga/src/front/glsl/functions.rs
+++ b/naga/src/front/glsl/functions.rs
@@ -1031,7 +1031,9 @@ impl Frontend {
         result: Option<FunctionResult>,
         meta: Span,
     ) {
-        ensure_block_returns(&mut ctx.body);
+        if result.is_some() {
+            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 d25eb362c1..f1e71e2c09 100644
--- a/naga/src/front/wgsl/lower/mod.rs
+++ b/naga/src/front/wgsl/lower/mod.rs
@@ -1307,7 +1307,9 @@ 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)?;
-        ensure_block_returns(&mut body);
+        if function.result.is_some() {
+            ensure_block_returns(&mut body);
+        }
 
         function.body = body;
         function.named_expressions = named_expressions
diff --git a/naga/src/proc/terminator.rs b/naga/src/proc/terminator.rs
index 19c37294ec..f22e61e6a6 100644
--- a/naga/src/proc/terminator.rs
+++ b/naga/src/proc/terminator.rs
@@ -28,9 +28,10 @@ pub fn ensure_block_returns(block: &mut crate::Block) {
                 }
             }
         }
-        Some(&mut (S::Emit(_) | S::Break | S::Continue | S::Return { .. } | S::Kill)) => (),
+        Some(&mut (S::Break | S::Continue | S::Return { .. } | S::Kill)) => (),
         Some(
-            &mut (S::Loop { .. }
+            &mut (S::Emit(_)
+            | S::Loop { .. }
             | S::Store { .. }
             | S::ImageStore { .. }
             | S::Call { .. }


Test patch:
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 004720bfc3..d269990a66 100644
--- a/naga/tests/out/glsl/access.assign_through_ptr.Compute.glsl
+++ b/naga/tests/out/glsl/access.assign_through_ptr.Compute.glsl
@@ -44,12 +44,10 @@ 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) {
@@ -59,7 +57,6 @@ 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]) {
@@ -69,7 +66,6 @@ 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) {
@@ -114,6 +110,5 @@ 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 b0224010aa..2835553237 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,12 +44,10 @@ 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) {
@@ -59,7 +57,6 @@ 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]) {
@@ -69,7 +66,6 @@ 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) {
@@ -116,6 +112,5 @@ 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 308cf195c1..48333b727a 100644
--- a/naga/tests/out/glsl/access.foo_frag.Fragment.glsl
+++ b/naga/tests/out/glsl/access.foo_frag.Fragment.glsl
@@ -54,12 +54,10 @@ 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) {
@@ -69,7 +67,6 @@ 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]) {
@@ -79,7 +76,6 @@ 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 c710043d23..870e603bc3 100644
--- a/naga/tests/out/glsl/access.foo_vert.Vertex.glsl
+++ b/naga/tests/out/glsl/access.foo_vert.Vertex.glsl
@@ -77,7 +77,6 @@ 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() {
@@ -113,7 +112,6 @@ 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) {
@@ -127,12 +125,10 @@ 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) {
@@ -142,7 +138,6 @@ 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]) {
@@ -152,7 +147,6 @@ 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/atomicOps.cs_main.Compute.glsl b/naga/tests/out/glsl/atomicOps.cs_main.Compute.glsl
index b69c5107ce..64a6a9e931 100644
--- a/naga/tests/out/glsl/atomicOps.cs_main.Compute.glsl
+++ b/naga/tests/out/glsl/atomicOps.cs_main.Compute.glsl
@@ -127,6 +127,5 @@ 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 4defe425f9..993bdfea67 100644
--- a/naga/tests/out/glsl/atomicTexture.cs_main.Compute.glsl
+++ b/naga/tests/out/glsl/atomicTexture.cs_main.Compute.glsl
@@ -22,6 +22,5 @@ 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 57e7e221b7..90d42e9e96 100644
--- a/naga/tests/out/glsl/bitcast.main.Compute.glsl
+++ b/naga/tests/out/glsl/bitcast.main.Compute.glsl
@@ -34,6 +34,5 @@ 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 f4b5c0f481..8558e6b402 100644
--- a/naga/tests/out/glsl/bits.main.Compute.glsl
+++ b/naga/tests/out/glsl/bits.main.Compute.glsl
@@ -129,6 +129,5 @@ 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 c42358bfef..b9db93cdb2 100644
--- a/naga/tests/out/glsl/boids.main.Compute.glsl
+++ b/naga/tests/out/glsl/boids.main.Compute.glsl
@@ -150,6 +150,5 @@ 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 f582e85ec7..af6bad24bd 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,27 +58,22 @@ 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 b2096add79..c7365c9046 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,27 +52,22 @@ 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 ba96de4879..1bcfe54006 100644
--- a/naga/tests/out/glsl/break-if.main.Compute.glsl
+++ b/naga/tests/out/glsl/break-if.main.Compute.glsl
@@ -16,7 +16,6 @@ void breakIfEmpty() {
         }
         loop_init = false;
     }
-    return;
 }
 
 void breakIfEmptyBody(bool a) {
@@ -35,7 +34,6 @@ void breakIfEmptyBody(bool a) {
         }
         loop_init_1 = false;
     }
-    return;
 }
 
 void breakIf(bool a_1) {
@@ -54,7 +52,6 @@ void breakIf(bool a_1) {
         bool _e2 = d;
         e = (a_1 != _e2);
     }
-    return;
 }
 
 void breakIfSeparateVariable() {
@@ -71,10 +68,8 @@ 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 b095345de9..7e6e9ad993 100644
--- a/naga/tests/out/glsl/const-exprs.main.Compute.glsl
+++ b/naga/tests/out/glsl/const-exprs.main.Compute.glsl
@@ -48,7 +48,6 @@ void non_constant_initializers() {
     int _e10 = y;
     int _e11 = z;
     out_3 = ivec4(_e8, _e9, _e10, _e11);
-    return;
 }
 
 void splat_of_constant() {
@@ -88,6 +87,5 @@ void main() {
     splat_of_constant();
     compose_of_constant();
     compose_of_splat();
-    return;
 }
 
diff --git a/naga/tests/out/glsl/control-flow.main.Compute.glsl b/naga/tests/out/glsl/control-flow.main.Compute.glsl
index 391fca84f4..4e9f63c405 100644
--- a/naga/tests/out/glsl/control-flow.main.Compute.glsl
+++ b/naga/tests/out/glsl/control-flow.main.Compute.glsl
@@ -35,7 +35,6 @@ void loop_switch_continue(int x) {
             }
         }
     }
-    return;
 }
 
 void loop_switch_continue_nesting(int x_1, int y, int z) {
@@ -93,7 +92,6 @@ 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) {
@@ -139,7 +137,6 @@ void loop_switch_omit_continue_variable_checks(int x_2, int y_1, int z_1, int w)
             }
         }
     }
-    return;
 }
 
 void main() {
@@ -193,18 +190,18 @@ void main() {
         }
         case 2: {
             pos = 1;
-            return;
+            break;
         }
         case 3: {
             pos = 2;
-            return;
+            break;
         }
         case 4: {
-            return;
+            break;
         }
         default: {
             pos = 3;
-            return;
+            break;
         }
     }
 }
diff --git a/naga/tests/out/glsl/empty.main.Compute.glsl b/naga/tests/out/glsl/empty.main.Compute.glsl
index f6b9108cc6..6e113755b9 100644
--- a/naga/tests/out/glsl/empty.main.Compute.glsl
+++ b/naga/tests/out/glsl/empty.main.Compute.glsl
@@ -7,6 +7,5 @@ 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 055965d153..42ed66e573 100644
--- a/naga/tests/out/glsl/f64.main.Compute.glsl
+++ b/naga/tests/out/glsl/f64.main.Compute.glsl
@@ -14,6 +14,5 @@ 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 9dd084c32a..bd7dea43bd 100644
--- a/naga/tests/out/glsl/functions-webgl.main.Fragment.glsl
+++ b/naga/tests/out/glsl/functions-webgl.main.Fragment.glsl
@@ -13,6 +13,5 @@ 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 b0f23125f7..ce5e351429 100644
--- a/naga/tests/out/glsl/functions.main.Compute.glsl
+++ b/naga/tests/out/glsl/functions.main.Compute.glsl
@@ -29,6 +29,5 @@ 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 b7ef8bd295..a72ebabdcc 100644
--- a/naga/tests/out/glsl/globals.main.Compute.glsl
+++ b/naga/tests/out/glsl/globals.main.Compute.glsl
@@ -31,7 +31,6 @@ 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_() {
@@ -78,6 +77,5 @@ 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 13a4b24d6e..2a38f35f37 100644
--- a/naga/tests/out/glsl/image.main.Compute.glsl
+++ b/naga/tests/out/glsl/image.main.Compute.glsl
@@ -33,6 +33,5 @@ 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 ac7fad324f..047d0b56b3 100644
--- a/naga/tests/out/glsl/interpolate_compat.frag_main.Fragment.glsl
+++ b/naga/tests/out/glsl/interpolate_compat.frag_main.Fragment.glsl
@@ -25,6 +25,5 @@ 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 9936a28ad3..724ff9bb3d 100644
--- a/naga/tests/out/glsl/invariant.fs.Fragment.glsl
+++ b/naga/tests/out/glsl/invariant.fs.Fragment.glsl
@@ -6,6 +6,5 @@ precision highp int;
 
 void main() {
     vec4 position = gl_FragCoord;
-    return;
 }
 
diff --git a/naga/tests/out/glsl/multiview.main.Fragment.glsl b/naga/tests/out/glsl/multiview.main.Fragment.glsl
index 466aea062f..775d5609bc 100644
--- a/naga/tests/out/glsl/multiview.main.Fragment.glsl
+++ b/naga/tests/out/glsl/multiview.main.Fragment.glsl
@@ -7,6 +7,5 @@ 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 30515289c9..50e785a434 100644
--- a/naga/tests/out/glsl/multiview_webgl.main.Fragment.glsl
+++ b/naga/tests/out/glsl/multiview_webgl.main.Fragment.glsl
@@ -7,6 +7,5 @@ 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 006bce205e..e11c26156c 100644
--- a/naga/tests/out/glsl/operators.main.Compute.glsl
+++ b/naga/tests/out/glsl/operators.main.Compute.glsl
@@ -232,7 +232,6 @@ void assignment() {
     vec0_[1] = (_e37 + 1);
     int _e41 = vec0_[1];
     vec0_[1] = (_e41 - 1);
-    return;
 }
 
 void negation_avoids_prefix_decrement() {
@@ -256,6 +255,5 @@ 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 684e7cca58..b1c631cec0 100644
--- a/naga/tests/out/glsl/overrides.main.Compute.glsl
+++ b/naga/tests/out/glsl/overrides.main.Compute.glsl
@@ -27,6 +27,5 @@ void main() {
     float _e9 = gain_x_10_;
     gain_x_100_ = (_e9 * 10.0);
     store_override = gain;
-    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 869b7ca418..03d40c7a6c 100644
--- a/naga/tests/out/glsl/separate-entry-points.compute.Compute.glsl
+++ b/naga/tests/out/glsl/separate-entry-points.compute.Compute.glsl
@@ -11,11 +11,9 @@ void barriers() {
     barrier();
     memoryBarrierShared();
     barrier();
-    return;
 }
 
 void main() {
     barriers();
-    return;
 }
 
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 7fe97f2982..8e682aed72 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,6 +25,5 @@ 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 c7bbcf7b70..6e8a367c89 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,6 +25,5 @@ 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 05ab403565..ef3eb5ac95 100644
--- a/naga/tests/out/glsl/subgroup-operations.main.Compute.glsl
+++ b/naga/tests/out/glsl/subgroup-operations.main.Compute.glsl
@@ -40,6 +40,5 @@ 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 25b258987c..1651a5e805 100644
--- a/naga/tests/out/glsl/variations.main.Fragment.glsl
+++ b/naga/tests/out/glsl/variations.main.Fragment.glsl
@@ -10,7 +10,6 @@ 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 6315309c99..ee50e2e1b6 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,9 +25,6 @@ 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 de136c1109..783afdaa52 100644
--- a/naga/tests/out/glsl/workgroup-var-init.main.Compute.glsl
+++ b/naga/tests/out/glsl/workgroup-var-init.main.Compute.glsl
@@ -23,6 +23,5 @@ void main() {
     barrier();
     uint _e3[512] = w_mem.arr;
     _group_0_binding_0_cs = _e3;
-    return;
 }
 
diff --git a/naga/tests/out/hlsl/access.hlsl b/naga/tests/out/hlsl/access.hlsl
index dcb62ae2af..9452f354fd 100644
--- a/naga/tests/out/hlsl/access.hlsl
+++ b/naga/tests/out/hlsl/access.hlsl
@@ -166,7 +166,6 @@ void test_matrix_within_struct_accesses()
     int _e85 = idx;
     int _e87 = idx;
     SetMatScalarmOnBaz(t, 40.0, _e85, _e87);
-    return;
 }
 
 MatCx2InArray ConstructMatCx2InArray(float4x2 arg0[2]) {
@@ -215,7 +214,6 @@ 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)
@@ -232,7 +230,6 @@ 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];
@@ -244,7 +241,6 @@ 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)
@@ -256,7 +252,6 @@ 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])
@@ -268,7 +263,6 @@ 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];
@@ -409,7 +403,6 @@ void assign_through_ptr()
 
     assign_through_ptr_fn(val);
     assign_array_through_ptr_fn(arr);
-    return;
 }
 
 [numthreads(1, 1, 1)]
@@ -422,5 +415,4 @@ 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/atomicOps-int64-min-max.hlsl b/naga/tests/out/hlsl/atomicOps-int64-min-max.hlsl
index 989a52b78b..81653010b2 100644
--- a/naga/tests/out/hlsl/atomicOps-int64-min-max.hlsl
+++ b/naga/tests/out/hlsl/atomicOps-int64-min-max.hlsl
@@ -31,5 +31,4 @@ 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 ea88f81753..34af783426 100644
--- a/naga/tests/out/hlsl/atomicOps-int64.hlsl
+++ b/naga/tests/out/hlsl/atomicOps-int64.hlsl
@@ -114,5 +114,4 @@ 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 640972a2fa..29588c4353 100644
--- a/naga/tests/out/hlsl/atomicOps.hlsl
+++ b/naga/tests/out/hlsl/atomicOps.hlsl
@@ -107,5 +107,4 @@ 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 056489c790..19f087b34a 100644
--- a/naga/tests/out/hlsl/atomicTexture-int64.hlsl
+++ b/naga/tests/out/hlsl/atomicTexture-int64.hlsl
@@ -13,5 +13,4 @@ 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 241cdab678..40f5789585 100644
--- a/naga/tests/out/hlsl/atomicTexture.hlsl
+++ b/naga/tests/out/hlsl/atomicTexture.hlsl
@@ -23,5 +23,4 @@ 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 b21408dda5..4e33e7aafa 100644
--- a/naga/tests/out/hlsl/bitcast.hlsl
+++ b/naga/tests/out/hlsl/bitcast.hlsl
@@ -29,5 +29,4 @@ 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 1987e15514..75adce1636 100644
--- a/naga/tests/out/hlsl/bits.hlsl
+++ b/naga/tests/out/hlsl/bits.hlsl
@@ -300,5 +300,4 @@ 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 22e9c6cefd..0f881f949a 100644
--- a/naga/tests/out/hlsl/boids.hlsl
+++ b/naga/tests/out/hlsl/boids.hlsl
@@ -140,5 +140,4 @@ 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 9e92c1ed6d..969300ab8e 100644
--- a/naga/tests/out/hlsl/bounds-check-dynamic-buffer.hlsl
+++ b/naga/tests/out/hlsl/bounds-check-dynamic-buffer.hlsl
@@ -35,5 +35,4 @@ 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 63a0185583..ee8719776d 100644
--- a/naga/tests/out/hlsl/break-if.hlsl
+++ b/naga/tests/out/hlsl/break-if.hlsl
@@ -9,7 +9,6 @@ void breakIfEmpty()
         }
         loop_init = false;
     }
-    return;
 }
 
 void breakIfEmptyBody(bool a)
@@ -30,7 +29,6 @@ void breakIfEmptyBody(bool a)
         }
         loop_init_1 = false;
     }
-    return;
 }
 
 void breakIf(bool a_1)
@@ -51,7 +49,6 @@ void breakIf(bool a_1)
         bool _e2 = d;
         e = (a_1 != _e2);
     }
-    return;
 }
 
 void breakIfSeparateVariable()
@@ -70,11 +67,9 @@ 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 b00586aa4c..bfea10d3c7 100644
--- a/naga/tests/out/hlsl/collatz.hlsl
+++ b/naga/tests/out/hlsl/collatz.hlsl
@@ -35,5 +35,4 @@ 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 29bec5f17a..5f049f1b2f 100644
--- a/naga/tests/out/hlsl/const-exprs.hlsl
+++ b/naga/tests/out/hlsl/const-exprs.hlsl
@@ -48,7 +48,6 @@ void non_constant_initializers()
     int _e10 = y;
     int _e11 = z;
     out_3 = int4(_e8, _e9, _e10, _e11);
-    return;
 }
 
 void splat_of_constant()
@@ -97,5 +96,4 @@ void main()
     splat_of_constant();
     compose_of_constant();
     compose_of_splat();
-    return;
 }
diff --git a/naga/tests/out/hlsl/control-flow.hlsl b/naga/tests/out/hlsl/control-flow.hlsl
index 2438858a8a..d7333eadeb 100644
--- a/naga/tests/out/hlsl/control-flow.hlsl
+++ b/naga/tests/out/hlsl/control-flow.hlsl
@@ -35,7 +35,6 @@ void loop_switch_continue(int x)
             continue;
         }
     }
-    return;
 }
 
 void loop_switch_continue_nesting(int x_1, int y, int z)
@@ -108,7 +107,6 @@ 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)
@@ -165,7 +163,6 @@ void loop_switch_omit_continue_variable_checks(int x_2, int y_1, int z_1, int w)
             continue;
         }
     }
-    return;
 }
 
 [numthreads(1, 1, 1)]
@@ -219,18 +216,18 @@ void main(uint3 global_id : SV_DispatchThreadID)
         }
         case 2: {
             pos = 1;
-            return;
+            break;
         }
         case 3: {
             pos = 2;
-            return;
+            break;
         }
         case 4: {
-            return;
+            break;
         }
         default: {
             pos = 3;
-            return;
+            break;
         }
     }
 }
diff --git a/naga/tests/out/hlsl/empty.hlsl b/naga/tests/out/hlsl/empty.hlsl
index e79d36b9c8..e9f2b4acd0 100644
--- a/naga/tests/out/hlsl/empty.hlsl
+++ b/naga/tests/out/hlsl/empty.hlsl
@@ -1,5 +1,4 @@
 [numthreads(1, 1, 1)]
 void main()
 {
-    return;
 }
diff --git a/naga/tests/out/hlsl/f64.hlsl b/naga/tests/out/hlsl/f64.hlsl
index 7b0fb9f67e..57a74cadc9 100644
--- a/naga/tests/out/hlsl/f64.hlsl
+++ b/naga/tests/out/hlsl/f64.hlsl
@@ -15,5 +15,4 @@ 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 6d7e210307..b765e59754 100644
--- a/naga/tests/out/hlsl/functions.hlsl
+++ b/naga/tests/out/hlsl/functions.hlsl
@@ -23,5 +23,4 @@ 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 d6d8eb4107..2f713eeb4b 100644
--- a/naga/tests/out/hlsl/globals.hlsl
+++ b/naga/tests/out/hlsl/globals.hlsl
@@ -68,7 +68,6 @@ cbuffer global_nested_arrays_of_matrices_4x2_ : register(b7) { __mat4x2 global_n
 
 void test_msl_packed_vec3_as_arg(float3 arg)
 {
-    return;
 }
 
 float3x3 ZeroValuefloat3x3() {
@@ -137,5 +136,4 @@ 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 30bfd86aeb..d749d1a688 100644
--- a/naga/tests/out/hlsl/image.hlsl
+++ b/naga/tests/out/hlsl/image.hlsl
@@ -50,7 +50,6 @@ 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 26441380b1..9d60a8acba 100644
--- a/naga/tests/out/hlsl/int64.hlsl
+++ b/naga/tests/out/hlsl/int64.hlsl
@@ -230,5 +230,4 @@ 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 6187ca0974..30e2c45b75 100644
--- a/naga/tests/out/hlsl/interface.hlsl
+++ b/naga/tests/out/hlsl/interface.hlsl
@@ -82,7 +82,6 @@ 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 ebb6e83477..c2ba2313ee 100644
--- a/naga/tests/out/hlsl/interpolate.hlsl
+++ b/naga/tests/out/hlsl/interpolate.hlsl
@@ -68,5 +68,4 @@ 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 85f1bb001c..7caf19904f 100644
--- a/naga/tests/out/hlsl/interpolate_compat.hlsl
+++ b/naga/tests/out/hlsl/interpolate_compat.hlsl
@@ -64,5 +64,4 @@ 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/operators.hlsl b/naga/tests/out/hlsl/operators.hlsl
index 7d9dc8f401..94e25661f6 100644
--- a/naga/tests/out/hlsl/operators.hlsl
+++ b/naga/tests/out/hlsl/operators.hlsl
@@ -251,7 +251,6 @@ void assignment()
     vec0_[1] = (_e37 + 1);
     int _e41 = vec0_[1];
     vec0_[1] = (_e41 - 1);
-    return;
 }
 
 void negation_avoids_prefix_decrement()
@@ -277,5 +276,4 @@ 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 5cb0e75687..0b1f7852eb 100644
--- a/naga/tests/out/hlsl/overrides.hlsl
+++ b/naga/tests/out/hlsl/overrides.hlsl
@@ -21,5 +21,4 @@ void main()
     float _e9 = gain_x_10_;
     gain_x_100_ = (_e9 * 10.0);
     store_override = gain;
-    return;
 }
diff --git a/naga/tests/out/hlsl/ray-query.hlsl b/naga/tests/out/hlsl/ray-query.hlsl
index 9a0a2da1ce..f4aa54b2e1 100644
--- a/naga/tests/out/hlsl/ray-query.hlsl
+++ b/naga/tests/out/hlsl/ray-query.hlsl
@@ -114,7 +114,6 @@ 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) {
@@ -148,5 +147,4 @@ 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 59f046ded6..cc0ac76b0c 100644
--- a/naga/tests/out/hlsl/struct-layout.hlsl
+++ b/naga/tests/out/hlsl/struct-layout.hlsl
@@ -52,7 +52,6 @@ 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
@@ -83,5 +82,4 @@ 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 839b1fa6b2..ab92f3cc36 100644
--- a/naga/tests/out/hlsl/subgroup-operations.hlsl
+++ b/naga/tests/out/hlsl/subgroup-operations.hlsl
@@ -34,5 +34,4 @@ 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 a8dba16b9a..b8bfa44930 100644
--- a/naga/tests/out/hlsl/workgroup-uniform-load.hlsl
+++ b/naga/tests/out/hlsl/workgroup-uniform-load.hlsl
@@ -14,8 +14,5 @@ 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 49b4fe621a..917f0b8845 100644
--- a/naga/tests/out/hlsl/workgroup-var-init.hlsl
+++ b/naga/tests/out/hlsl/workgroup-var-init.hlsl
@@ -530,5 +530,4 @@ 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 53ba1a3c9b..8e05e2936c 100644
--- a/naga/tests/out/ir/access.compact.ron
+++ b/naga/tests/out/ir/access.compact.ron
@@ -982,9 +982,6 @@
                     pointer: 88,
                     value: 89,
                 ),
-                Return(
-                    value: None,
-                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -1587,9 +1584,6 @@
                     pointer: 103,
                     value: 104,
                 ),
-                Return(
-                    value: None,
-                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -1693,9 +1687,6 @@
                     pointer: 0,
                     value: 1,
                 ),
-                Return(
-                    value: None,
-                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -1754,9 +1745,6 @@
                     pointer: 0,
                     value: 5,
                 ),
-                Return(
-                    value: None,
-                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -1829,9 +1817,6 @@
                     pointer: 1,
                     value: 2,
                 ),
-                Return(
-                    value: None,
-                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -1904,9 +1889,6 @@
                     pointer: 1,
                     value: 2,
                 ),
-                Return(
-                    value: None,
-                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -2779,9 +2761,6 @@
                         ],
                         result: None,
                     ),
-                    Return(
-                        value: None,
-                    ),
                 ],
                 diagnostic_filter_leaf: None,
             ),
@@ -2844,9 +2823,6 @@
                         ],
                         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 53ba1a3c9b..8e05e2936c 100644
--- a/naga/tests/out/ir/access.ron
+++ b/naga/tests/out/ir/access.ron
@@ -982,9 +982,6 @@
                     pointer: 88,
                     value: 89,
                 ),
-                Return(
-                    value: None,
-                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -1587,9 +1584,6 @@
                     pointer: 103,
                     value: 104,
                 ),
-                Return(
-                    value: None,
-                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -1693,9 +1687,6 @@
                     pointer: 0,
                     value: 1,
                 ),
-                Return(
-                    value: None,
-                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -1754,9 +1745,6 @@
                     pointer: 0,
                     value: 5,
                 ),
-                Return(
-                    value: None,
-                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -1829,9 +1817,6 @@
                     pointer: 1,
                     value: 2,
                 ),
-                Return(
-                    value: None,
-                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -1904,9 +1889,6 @@
                     pointer: 1,
                     value: 2,
                 ),
-                Return(
-                    value: None,
-                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -2779,9 +2761,6 @@
                         ],
                         result: None,
                     ),
-                    Return(
-                        value: None,
-                    ),
                 ],
                 diagnostic_filter_leaf: None,
             ),
@@ -2844,9 +2823,6 @@
                         ],
                         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 6a7aebe544..d4612a7599 100644
--- a/naga/tests/out/ir/collatz.compact.ron
+++ b/naga/tests/out/ir/collatz.compact.ron
@@ -325,9 +325,6 @@
                         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 6a7aebe544..d4612a7599 100644
--- a/naga/tests/out/ir/collatz.ron
+++ b/naga/tests/out/ir/collatz.ron
@@ -325,9 +325,6 @@
                         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 03c540b601..cc2f6aa035 100644
--- a/naga/tests/out/ir/const_assert.compact.ron
+++ b/naga/tests/out/ir/const_assert.compact.ron
@@ -43,11 +43,7 @@
             named_expressions: {
                 0: "z",
             },
-            body: [
-                Return(
-                    value: None,
-                ),
-            ],
+            body: [],
             diagnostic_filter_leaf: None,
         ),
     ],
diff --git a/naga/tests/out/ir/const_assert.ron b/naga/tests/out/ir/const_assert.ron
index 03c540b601..cc2f6aa035 100644
--- a/naga/tests/out/ir/const_assert.ron
+++ b/naga/tests/out/ir/const_assert.ron
@@ -43,11 +43,7 @@
             named_expressions: {
                 0: "z",
             },
-            body: [
-                Return(
-                    value: None,
-                ),
-            ],
+            body: [],
             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 315c8b3423..df8f1c9161 100644
--- a/naga/tests/out/ir/diagnostic-filter.compact.ron
+++ b/naga/tests/out/ir/diagnostic-filter.compact.ron
@@ -17,11 +17,7 @@
             local_variables: [],
             expressions: [],
             named_expressions: {},
-            body: [
-                Return(
-                    value: None,
-                ),
-            ],
+            body: [],
             diagnostic_filter_leaf: Some(0),
         ),
         (
@@ -31,11 +27,7 @@
             local_variables: [],
             expressions: [],
             named_expressions: {},
-            body: [
-                Return(
-                    value: None,
-                ),
-            ],
+            body: [],
             diagnostic_filter_leaf: Some(1),
         ),
     ],
diff --git a/naga/tests/out/ir/diagnostic-filter.ron b/naga/tests/out/ir/diagnostic-filter.ron
index 315c8b3423..df8f1c9161 100644
--- a/naga/tests/out/ir/diagnostic-filter.ron
+++ b/naga/tests/out/ir/diagnostic-filter.ron
@@ -17,11 +17,7 @@
             local_variables: [],
             expressions: [],
             named_expressions: {},
-            body: [
-                Return(
-                    value: None,
-                ),
-            ],
+            body: [],
             diagnostic_filter_leaf: Some(0),
         ),
         (
@@ -31,11 +27,7 @@
             local_variables: [],
             expressions: [],
             named_expressions: {},
-            body: [
-                Return(
-                    value: None,
-                ),
-            ],
+            body: [],
             diagnostic_filter_leaf: Some(1),
         ),
     ],
diff --git a/naga/tests/out/ir/must-use.compact.ron b/naga/tests/out/ir/must-use.compact.ron
index 3d51cb0c95..c29feff15e 100644
--- a/naga/tests/out/ir/must-use.compact.ron
+++ b/naga/tests/out/ir/must-use.compact.ron
@@ -146,9 +146,6 @@
                     arguments: [],
                     result: Some(0),
                 ),
-                Return(
-                    value: None,
-                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -167,11 +164,7 @@
                 local_variables: [],
                 expressions: [],
                 named_expressions: {},
-                body: [
-                    Return(
-                        value: None,
-                    ),
-                ],
+                body: [],
                 diagnostic_filter_leaf: None,
             ),
         ),
diff --git a/naga/tests/out/ir/must-use.ron b/naga/tests/out/ir/must-use.ron
index 3d51cb0c95..c29feff15e 100644
--- a/naga/tests/out/ir/must-use.ron
+++ b/naga/tests/out/ir/must-use.ron
@@ -146,9 +146,6 @@
                     arguments: [],
                     result: Some(0),
                 ),
-                Return(
-                    value: None,
-                ),
             ],
             diagnostic_filter_leaf: None,
         ),
@@ -167,11 +164,7 @@
                 local_variables: [],
                 expressions: [],
                 named_expressions: {},
-                body: [
-                    Return(
-                        value: None,
-                    ),
-                ],
+                body: [],
                 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 56be2f8ab6..6dd480e0b0 100644
--- a/naga/tests/out/ir/overrides-atomicCompareExchangeWeak.compact.ron
+++ b/naga/tests/out/ir/overrides-atomicCompareExchangeWeak.compact.ron
@@ -119,9 +119,6 @@
                         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 56be2f8ab6..6dd480e0b0 100644
--- a/naga/tests/out/ir/overrides-atomicCompareExchangeWeak.ron
+++ b/naga/tests/out/ir/overrides-atomicCompareExchangeWeak.ron
@@ -119,9 +119,6 @@
                         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 10cad83538..5d4964ae67 100644
--- a/naga/tests/out/ir/overrides-ray-query.compact.ron
+++ b/naga/tests/out/ir/overrides-ray-query.compact.ron
@@ -250,9 +250,6 @@
                         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 10cad83538..5d4964ae67 100644
--- a/naga/tests/out/ir/overrides-ray-query.ron
+++ b/naga/tests/out/ir/overrides-ray-query.ron
@@ -250,9 +250,6 @@
                         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 00c57fa434..8fd68b84f6 100644
--- a/naga/tests/out/ir/overrides.compact.ron
+++ b/naga/tests/out/ir/overrides.compact.ron
@@ -204,9 +204,6 @@
                         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 00c57fa434..8fd68b84f6 100644
--- a/naga/tests/out/ir/overrides.ron
+++ b/naga/tests/out/ir/overrides.ron
@@ -204,9 +204,6 @@
                         pointer: 11,
                         value: 12,
                     ),
-                    Return(
-                        value: None,
-                    ),
                 ],
                 diagnostic_filter_leaf: None,
             ),
diff --git a/naga/tests/out/msl/abstract-types-function-calls.msl b/naga/tests/out/msl/abstract-types-function-calls.msl
index bf62d87a54..c3c6d812c8 100644
--- a/naga/tests/out/msl/abstract-types-function-calls.msl
+++ b/naga/tests/out/msl/abstract-types-function-calls.msl
@@ -17,68 +17,57 @@ 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_(
@@ -99,5 +88,4 @@ 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 54fd234d29..633164432b 100644
--- a/naga/tests/out/msl/abstract-types-operators.msl
+++ b/naga/tests/out/msl/abstract-types-operators.msl
@@ -86,12 +86,10 @@ void runtime_values(
     uint _e52 = u;
     uint _e53 = u;
     plus_u_u_u = _e52 + _e53;
-    return;
 }
 
 void wgpu_4445_(
 ) {
-    return;
 }
 
 void wgpu_4435_(
diff --git a/naga/tests/out/msl/abstract-types-var.msl b/naga/tests/out/msl/abstract-types-var.msl
index 1c9d6fcb25..ed84db393d 100644
--- a/naga/tests/out/msl/abstract-types-var.msl
+++ b/naga/tests/out/msl/abstract-types-var.msl
@@ -87,7 +87,6 @@ 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(
@@ -195,5 +194,4 @@ 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 eb0808f80e..e5a477517e 100644
--- a/naga/tests/out/msl/access.msl
+++ b/naga/tests/out/msl/access.msl
@@ -110,7 +110,6 @@ 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(
@@ -148,7 +147,6 @@ 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(
@@ -168,14 +166,12 @@ 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(
@@ -189,7 +185,6 @@ void assign_to_arg_ptr_member(
     thread AssignToMember& p_2
 ) {
     p_2.x = 10u;
-    return;
 }
 
 uint fetch_arg_ptr_array_element(
@@ -203,7 +198,6 @@ void assign_to_arg_ptr_array_element(
     thread type_25& p_4
 ) {
     p_4.inner[1] = 10u;
-    return;
 }
 
 bool index_ptr(
@@ -304,7 +298,6 @@ 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;
 }
 
 
@@ -316,5 +309,4 @@ 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/atomicCompareExchange.msl b/naga/tests/out/msl/atomicCompareExchange.msl
index 6655fad7e2..7ec7eb272e 100644
--- a/naga/tests/out/msl/atomicCompareExchange.msl
+++ b/naga/tests/out/msl/atomicCompareExchange.msl
@@ -114,7 +114,6 @@ kernel void test_atomic_compare_exchange_i32_(
         }
         LOOP_IS_BOUNDED
     }
-    return;
 }
 
 
@@ -161,5 +160,4 @@ 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 40b5a7a14f..d062ff5852 100644
--- a/naga/tests/out/msl/atomicOps-float32.msl
+++ b/naga/tests/out/msl/atomicOps-float32.msl
@@ -39,5 +39,4 @@ 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 f69a2a49bd..998f54d275 100644
--- a/naga/tests/out/msl/atomicOps-int64-min-max.msl
+++ b/naga/tests/out/msl/atomicOps-int64-min-max.msl
@@ -34,5 +34,4 @@ 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 4732b4a32d..8d44e9d64c 100644
--- a/naga/tests/out/msl/atomicOps.msl
+++ b/naga/tests/out/msl/atomicOps.msl
@@ -122,5 +122,4 @@ 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 c00d8b7654..59704d05ff 100644
--- a/naga/tests/out/msl/atomicTexture-int64.msl
+++ b/naga/tests/out/msl/atomicTexture-int64.msl
@@ -14,5 +14,4 @@ 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 2cb17a081e..f61636f4e4 100644
--- a/naga/tests/out/msl/atomicTexture.msl
+++ b/naga/tests/out/msl/atomicTexture.msl
@@ -24,5 +24,4 @@ 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 20f4b850e3..80be97f2f3 100644
--- a/naga/tests/out/msl/bitcast.msl
+++ b/naga/tests/out/msl/bitcast.msl
@@ -34,5 +34,4 @@ 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 a0138862d5..2a4d51a391 100644
--- a/naga/tests/out/msl/bits.msl
+++ b/naga/tests/out/msl/bits.msl
@@ -129,5 +129,4 @@ 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 a4d709bd09..be74f8fb45 100644
--- a/naga/tests/out/msl/boids.msl
+++ b/naga/tests/out/msl/boids.msl
@@ -156,5 +156,4 @@ 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 138c0f6455..5da5ca711f 100644
--- a/naga/tests/out/msl/bounds-check-image-restrict.msl
+++ b/naga/tests/out/msl/bounds-check-image-restrict.msl
@@ -112,7 +112,6 @@ 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(
@@ -121,7 +120,6 @@ 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(
@@ -131,7 +129,6 @@ 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(
@@ -141,7 +138,6 @@ 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(
@@ -150,7 +146,6 @@ 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 f73b8e3e32..8eeb535535 100644
--- a/naga/tests/out/msl/bounds-check-image-rzsw.msl
+++ b/naga/tests/out/msl/bounds-check-image-rzsw.msl
@@ -111,7 +111,6 @@ 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,7 +119,6 @@ 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(
@@ -130,7 +128,6 @@ 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(
@@ -140,7 +137,6 @@ 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(
@@ -149,7 +145,6 @@ 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 a8bf371555..6202d878c1 100644
--- a/naga/tests/out/msl/bounds-check-restrict.msl
+++ b/naga/tests/out/msl/bounds-check-restrict.msl
@@ -99,7 +99,6 @@ void set_array(
     constant _mslBufferSizes& _buffer_sizes
 ) {
     globals.a.inner[metal::min(unsigned(i_7), 9u)] = v_1;
-    return;
 }
 
 void set_dynamic_array(
@@ -109,7 +108,6 @@ 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(
@@ -119,7 +117,6 @@ void set_vector(
     constant _mslBufferSizes& _buffer_sizes
 ) {
     globals.v[metal::min(unsigned(i_9), 3u)] = v_3;
-    return;
 }
 
 void set_matrix(
@@ -129,7 +126,6 @@ void set_matrix(
     constant _mslBufferSizes& _buffer_sizes
 ) {
     globals.m[metal::min(unsigned(i_10), 2u)] = v_4;
-    return;
 }
 
 void set_index_twice(
@@ -140,7 +136,6 @@ 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(
@@ -150,7 +145,6 @@ 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(
@@ -161,7 +155,6 @@ 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(
@@ -178,5 +171,4 @@ 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 8269d4bf70..14cac0d4ef 100644
--- a/naga/tests/out/msl/bounds-check-zero.msl
+++ b/naga/tests/out/msl/bounds-check-zero.msl
@@ -108,7 +108,6 @@ void set_array(
     if (uint(i_7) < 10) {
         globals.a.inner[i_7] = v_1;
     }
-    return;
 }
 
 void set_dynamic_array(
@@ -120,7 +119,6 @@ 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(
@@ -132,7 +130,6 @@ void set_vector(
     if (uint(i_9) < 4) {
         globals.v[i_9] = v_3;
     }
-    return;
 }
 
 void set_matrix(
@@ -144,7 +141,6 @@ void set_matrix(
     if (uint(i_10) < 3) {
         globals.m[i_10] = v_4;
     }
-    return;
 }
 
 void set_index_twice(
@@ -157,7 +153,6 @@ 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(
@@ -170,7 +165,6 @@ void set_expensive(
     if (uint(_e10) < 10) {
         globals.a.inner[_e10] = v_6;
     }
-    return;
 }
 
 void set_in_bounds(
@@ -181,7 +175,6 @@ 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(
@@ -200,5 +193,4 @@ 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 4d3397234b..5b9812436f 100644
--- a/naga/tests/out/msl/break-if.msl
+++ b/naga/tests/out/msl/break-if.msl
@@ -18,7 +18,6 @@ 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(
@@ -40,7 +39,6 @@ void breakIfEmptyBody(
         loop_init_1 = false;
         LOOP_IS_BOUNDED
     }
-    return;
 }
 
 void breakIf(
@@ -62,7 +60,6 @@ void breakIf(
         e = a_1 != _e2;
         LOOP_IS_BOUNDED
     }
-    return;
 }
 
 void breakIfSeparateVariable(
@@ -81,10 +78,8 @@ 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 e282d13abb..36dc36a25d 100644
--- a/naga/tests/out/msl/collatz.msl
+++ b/naga/tests/out/msl/collatz.msl
@@ -54,5 +54,4 @@ 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 7798ae62b3..d82f5b277f 100644
--- a/naga/tests/out/msl/const-exprs.msl
+++ b/naga/tests/out/msl/const-exprs.msl
@@ -50,7 +50,6 @@ void non_constant_initializers(
     int _e10 = y;
     int _e11 = z;
     out_3 = metal::int4(_e8, _e9, _e10, _e11);
-    return;
 }
 
 void splat_of_constant(
@@ -96,5 +95,4 @@ kernel void main_(
     splat_of_constant();
     compose_of_constant();
     compose_of_splat();
-    return;
 }
diff --git a/naga/tests/out/msl/control-flow.msl b/naga/tests/out/msl/control-flow.msl
index 1b35249f36..2bc471ae9e 100644
--- a/naga/tests/out/msl/control-flow.msl
+++ b/naga/tests/out/msl/control-flow.msl
@@ -43,7 +43,6 @@ 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(
@@ -103,7 +102,6 @@ void loop_switch_continue_nesting(
         }
         LOOP_IS_BOUNDED
     }
-    return;
 }
 
 void loop_switch_omit_continue_variable_checks(
@@ -156,7 +154,6 @@ void loop_switch_omit_continue_variable_checks(
         }
         LOOP_IS_BOUNDED
     }
-    return;
 }
 
 struct main_Input {
@@ -214,18 +211,18 @@ kernel void main_(
         }
         case 2: {
             pos = 1;
-            return;
+            break;
         }
         case 3: {
             pos = 2;
-            return;
+            break;
         }
         case 4: {
-            return;
+            break;
         }
         default: {
             pos = 3;
-            return;
+            break;
         }
     }
 }
diff --git a/naga/tests/out/msl/empty.msl b/naga/tests/out/msl/empty.msl
index 414cd22012..38cee7f85a 100644
--- a/naga/tests/out/msl/empty.msl
+++ b/naga/tests/out/msl/empty.msl
@@ -7,5 +7,4 @@ using metal::uint;
 
 kernel void main_(
 ) {
-    return;
 }
diff --git a/naga/tests/out/msl/functions.msl b/naga/tests/out/msl/functions.msl
index 42632f99be..7d799b52e3 100644
--- a/naga/tests/out/msl/functions.msl
+++ b/naga/tests/out/msl/functions.msl
@@ -31,5 +31,4 @@ 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 d2ed89ed46..4baa36d1f8 100644
--- a/naga/tests/out/msl/globals.msl
+++ b/naga/tests/out/msl/globals.msl
@@ -36,7 +36,6 @@ constant bool Foo_1 = true;
 void test_msl_packed_vec3_as_arg(
     metal::float3 arg
 ) {
-    return;
 }
 
 void test_msl_packed_vec3_(
@@ -96,5 +95,4 @@ 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 1a8567a0ca..11405d2db5 100644
--- a/naga/tests/out/msl/image.msl
+++ b/naga/tests/out/msl/image.msl
@@ -32,7 +32,6 @@ 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 369eb01c6f..d26c139c3c 100644
--- a/naga/tests/out/msl/int64.msl
+++ b/naga/tests/out/msl/int64.msl
@@ -209,5 +209,4 @@ 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 047873da9f..9dc33927e9 100644
--- a/naga/tests/out/msl/interface.msl
+++ b/naga/tests/out/msl/interface.msl
@@ -81,7 +81,6 @@ 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 c19005753f..fd3baea127 100644
--- a/naga/tests/out/msl/interpolate.msl
+++ b/naga/tests/out/msl/interpolate.msl
@@ -73,5 +73,4 @@ 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 e386c07db0..e861d71454 100644
--- a/naga/tests/out/msl/interpolate_compat.msl
+++ b/naga/tests/out/msl/interpolate_compat.msl
@@ -70,5 +70,4 @@ 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/operators.msl b/naga/tests/out/msl/operators.msl
index da483ed794..e6d747811c 100644
--- a/naga/tests/out/msl/operators.msl
+++ b/naga/tests/out/msl/operators.msl
@@ -242,7 +242,6 @@ 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(
@@ -270,5 +269,4 @@ 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 d87190c595..dba6d8f6cb 100644
--- a/naga/tests/out/msl/overrides-atomicCompareExchangeWeak.msl
+++ b/naga/tests/out/msl/overrides-atomicCompareExchangeWeak.msl
@@ -44,5 +44,4 @@ 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 3aa0ee0359..8dc7bb8371 100644
--- a/naga/tests/out/msl/overrides-ray-query.msl
+++ b/naga/tests/out/msl/overrides-ray-query.msl
@@ -43,5 +43,4 @@ 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 02770d613d..32b6fd5f2e 100644
--- a/naga/tests/out/msl/overrides.msl
+++ b/naga/tests/out/msl/overrides.msl
@@ -24,5 +24,4 @@ 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/ray-query.msl b/naga/tests/out/msl/ray-query.msl
index b8230fb2e8..8d9ebbf7c5 100644
--- a/naga/tests/out/msl/ray-query.msl
+++ b/naga/tests/out/msl/ray-query.msl
@@ -86,7 +86,6 @@ 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;
 }
 
 
@@ -105,5 +104,4 @@ 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 e1f266435b..644927b6b1 100644
--- a/naga/tests/out/msl/struct-layout.msl
+++ b/naga/tests/out/msl/struct-layout.msl
@@ -54,7 +54,6 @@ kernel void no_padding_comp(
     x = _e2;
     NoPadding _e4 = no_padding_storage;
     x = _e4;
-    return;
 }
 
 
@@ -99,5 +98,4 @@ 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 20a550a715..e32d570e31 100644
--- a/naga/tests/out/msl/subgroup-operations.msl
+++ b/naga/tests/out/msl/subgroup-operations.msl
@@ -40,5 +40,4 @@ 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 32495c198a..dfe129c37c 100644
--- a/naga/tests/out/msl/workgroup-uniform-load.msl
+++ b/naga/tests/out/msl/workgroup-uniform-load.msl
@@ -25,8 +25,5 @@ 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 991c8b014b..9d4e4d9d5c 100644
--- a/naga/tests/out/msl/workgroup-var-init.msl
+++ b/naga/tests/out/msl/workgroup-var-init.msl
@@ -36,5 +36,4 @@ 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 f3c3644b4f..8fdb0cc34c 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
-OpReturn
+OpBranch %128
 %131 = OpLabel
 OpStore %108 %75
-OpReturn
+OpBranch %128
 %132 = OpLabel
-OpReturn
+OpBranch %128
 %133 = OpLabel
 OpStore %108 %105
-OpReturn
+OpBranch %128
 %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 87f212a799..69c8cc842c 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: 40
+; Bound: 39
 OpCapability Shader
 %1 = OpExtInstImport "GLSL.std.450"
 OpMemoryModel Logical GLSL450
@@ -55,12 +55,10 @@ OpControlBarrier %29 %29 %30
 OpControlBarrier %29 %29 %30
 %36 = OpSGreaterThan  %24  %35 %16
 OpSelectionMerge %37 None
-OpBranchConditional %36 %38 %39
+OpBranchConditional %36 %38 %37
 %38 = OpLabel
 OpControlBarrier %29 %29 %30
-OpReturn
-%39 = OpLabel
-OpReturn
+OpBranch %37
 %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 a800b0f856..18256ae7d2 100644
--- a/naga/tests/out/wgsl/210-bevy-2d-shader.frag.wgsl
+++ b/naga/tests/out/wgsl/210-bevy-2d-shader.frag.wgsl
@@ -18,7 +18,6 @@ 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 52355cd404..4947398bf9 100644
--- a/naga/tests/out/wgsl/210-bevy-2d-shader.vert.wgsl
+++ b/naga/tests/out/wgsl/210-bevy-2d-shader.vert.wgsl
@@ -39,7 +39,6 @@ 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 223538e4e0..db49013db6 100644
--- a/naga/tests/out/wgsl/210-bevy-shader.vert.wgsl
+++ b/naga/tests/out/wgsl/210-bevy-shader.vert.wgsl
@@ -40,7 +40,6 @@ 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 71453bfe1c..6fdc5106e8 100644
--- a/naga/tests/out/wgsl/246-collatz.comp.wgsl
+++ b/naga/tests/out/wgsl/246-collatz.comp.wgsl
@@ -48,7 +48,6 @@ 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 c9ef2be74a..9b0a19ab19 100644
--- a/naga/tests/out/wgsl/277-casting.frag.wgsl
+++ b/naga/tests/out/wgsl/277-casting.frag.wgsl
@@ -1,7 +1,6 @@
 fn main_1() {
     var a: f32 = 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 6b3c4ac973..152fea63be 100644
--- a/naga/tests/out/wgsl/484-preprocessor-if.frag.wgsl
+++ b/naga/tests/out/wgsl/484-preprocessor-if.frag.wgsl
@@ -1,5 +1,4 @@
 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 e2672a9951..315ce3d03a 100644
--- a/naga/tests/out/wgsl/5246-dual-iteration.frag.wgsl
+++ b/naga/tests/out/wgsl/5246-dual-iteration.frag.wgsl
@@ -41,7 +41,6 @@ fn main_1() {
             x = (_e6 + 1i);
         }
     }
-    return;
 }
 
 @fragment 
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 8d6bc50af0..32f57cb1fb 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,7 +29,6 @@ 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 729e35a43f..6c0c0c720e 100644
--- a/naga/tests/out/wgsl/896-push-constant.frag.wgsl
+++ b/naga/tests/out/wgsl/896-push-constant.frag.wgsl
@@ -5,7 +5,6 @@ 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 2c468dd929..b774f07a6a 100644
--- a/naga/tests/out/wgsl/900-implicit-conversions.frag.wgsl
+++ b/naga/tests/out/wgsl/900-implicit-conversions.frag.wgsl
@@ -2,63 +2,54 @@ 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 ae45360074..91c3fcf9be 100644
--- a/naga/tests/out/wgsl/901-lhs-field-select.frag.wgsl
+++ b/naga/tests/out/wgsl/901-lhs-field-select.frag.wgsl
@@ -2,7 +2,6 @@ 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 ac59a0c3ef..284a570264 100644
--- a/naga/tests/out/wgsl/931-constant-emitting.frag.wgsl
+++ b/naga/tests/out/wgsl/931-constant-emitting.frag.wgsl
@@ -5,7 +5,6 @@ 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 bf6a281b49..74f4819de2 100644
--- a/naga/tests/out/wgsl/932-for-loop-if.frag.wgsl
+++ b/naga/tests/out/wgsl/932-for-loop-if.frag.wgsl
@@ -13,7 +13,6 @@ 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 419125bc6b..68a0f8c80f 100644
--- a/naga/tests/out/wgsl/abstract-types-function-calls.wgsl
+++ b/naga/tests/out/wgsl/abstract-types-function-calls.wgsl
@@ -1,45 +1,34 @@
 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() {
@@ -59,6 +48,5 @@ 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 caafb535ce..23e8046ec5 100644
--- a/naga/tests/out/wgsl/abstract-types-operators.wgsl
+++ b/naga/tests/out/wgsl/abstract-types-operators.wgsl
@@ -79,11 +79,9 @@ fn runtime_values() {
     let _e52 = u;
     let _e53 = u;
     plus_u_u_u = (_e52 + _e53);
-    return;
 }
 
 fn wgpu_4445_() {
-    return;
 }
 
 fn wgpu_4435_() {
diff --git a/naga/tests/out/wgsl/abstract-types-var.wgsl b/naga/tests/out/wgsl/abstract-types-var.wgsl
index cd35501732..ad3b08c912 100644
--- a/naga/tests/out/wgsl/abstract-types-var.wgsl
+++ b/naga/tests/out/wgsl/abstract-types-var.wgsl
@@ -111,7 +111,6 @@ 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() {
@@ -219,6 +218,5 @@ 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 fd4f742522..a43ecd8d2f 100644
--- a/naga/tests/out/wgsl/access.wgsl
+++ b/naga/tests/out/wgsl/access.wgsl
@@ -84,7 +84,6 @@ 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() {
@@ -121,7 +120,6 @@ 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 {
@@ -135,12 +133,10 @@ 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 {
@@ -150,7 +146,6 @@ 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 {
@@ -160,7 +155,6 @@ 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 {
@@ -243,7 +237,6 @@ fn assign_through_ptr() {
 
     assign_through_ptr_fn((&val));
     assign_array_through_ptr_fn((&arr));
-    return;
 }
 
 @compute @workgroup_size(1, 1, 1) 
@@ -255,5 +248,4 @@ 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/atomicCompareExchange-int64.wgsl b/naga/tests/out/wgsl/atomicCompareExchange-int64.wgsl
index 07cce9a89f..ea4c82bac2 100644
--- a/naga/tests/out/wgsl/atomicCompareExchange-int64.wgsl
+++ b/naga/tests/out/wgsl/atomicCompareExchange-int64.wgsl
@@ -44,7 +44,6 @@ fn test_atomic_compare_exchange_i64_() {
             i = (_e26 + 1u);
         }
     }
-    return;
 }
 
 @compute @workgroup_size(1, 1, 1) 
@@ -86,5 +85,4 @@ 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 17bd5144b1..19aeb44251 100644
--- a/naga/tests/out/wgsl/atomicCompareExchange.wgsl
+++ b/naga/tests/out/wgsl/atomicCompareExchange.wgsl
@@ -44,7 +44,6 @@ fn test_atomic_compare_exchange_i32_() {
             i = (_e27 + 1u);
         }
     }
-    return;
 }
 
 @compute @workgroup_size(1, 1, 1) 
@@ -86,5 +85,4 @@ 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 92e67b3428..0a4350daa2 100644
--- a/naga/tests/out/wgsl/atomicOps-float32.wgsl
+++ b/naga/tests/out/wgsl/atomicOps-float32.wgsl
@@ -31,5 +31,4 @@ 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 126758b0b5..88e7751d1c 100644
--- a/naga/tests/out/wgsl/atomicOps-int64-min-max.wgsl
+++ b/naga/tests/out/wgsl/atomicOps-int64-min-max.wgsl
@@ -27,5 +27,4 @@ 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 364108c646..c79c8e6cae 100644
--- a/naga/tests/out/wgsl/atomicOps-int64.wgsl
+++ b/naga/tests/out/wgsl/atomicOps-int64.wgsl
@@ -103,5 +103,4 @@ 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 be102e6833..696d6dd165 100644
--- a/naga/tests/out/wgsl/atomicOps.wgsl
+++ b/naga/tests/out/wgsl/atomicOps.wgsl
@@ -103,5 +103,4 @@ 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 52bbe76771..8cec19cbc7 100644
--- a/naga/tests/out/wgsl/atomicTexture-int64.wgsl
+++ b/naga/tests/out/wgsl/atomicTexture-int64.wgsl
@@ -6,5 +6,4 @@ 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 f5bd7db002..48c141fa2a 100644
--- a/naga/tests/out/wgsl/atomicTexture.wgsl
+++ b/naga/tests/out/wgsl/atomicTexture.wgsl
@@ -17,5 +17,4 @@ 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 78d3fbe606..c627967cc5 100644
--- a/naga/tests/out/wgsl/bevy-pbr.frag.wgsl
+++ b/naga/tests/out/wgsl/bevy-pbr.frag.wgsl
@@ -849,7 +849,6 @@ 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 1524b9e469..e8c85052c4 100644
--- a/naga/tests/out/wgsl/bevy-pbr.vert.wgsl
+++ b/naga/tests/out/wgsl/bevy-pbr.vert.wgsl
@@ -49,7 +49,6 @@ 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 1f36ce5e66..1c5e31502a 100644
--- a/naga/tests/out/wgsl/bitcast.wgsl
+++ b/naga/tests/out/wgsl/bitcast.wgsl
@@ -28,5 +28,4 @@ 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 0d23b1e782..23371c247e 100644
--- a/naga/tests/out/wgsl/bits.wgsl
+++ b/naga/tests/out/wgsl/bits.wgsl
@@ -123,5 +123,4 @@ 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 3513d1ba25..833398422e 100644
--- a/naga/tests/out/wgsl/bits_glsl.frag.wgsl
+++ b/naga/tests/out/wgsl/bits_glsl.frag.wgsl
@@ -102,7 +102,6 @@ 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 ee28ca6786..cb64da41fe 100644
--- a/naga/tests/out/wgsl/boids.wgsl
+++ b/naga/tests/out/wgsl/boids.wgsl
@@ -144,5 +144,4 @@ 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 4045a3de45..525c22ba6b 100644
--- a/naga/tests/out/wgsl/bool-select.frag.wgsl
+++ b/naga/tests/out/wgsl/bool-select.frag.wgsl
@@ -34,7 +34,6 @@ 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 c3d45a50ac..5356f3d866 100644
--- a/naga/tests/out/wgsl/break-if.wgsl
+++ b/naga/tests/out/wgsl/break-if.wgsl
@@ -4,7 +4,6 @@ fn breakIfEmpty() {
             break if true;
         }
     }
-    return;
 }
 
 fn breakIfEmptyBody(a: bool) {
@@ -20,7 +19,6 @@ fn breakIfEmptyBody(a: bool) {
             break if (a == _e5);
         }
     }
-    return;
 }
 
 fn breakIf(a_1: bool) {
@@ -36,7 +34,6 @@ fn breakIf(a_1: bool) {
             break if (a_1 == _e5);
         }
     }
-    return;
 }
 
 fn breakIfSeparateVariable() {
@@ -50,10 +47,8 @@ 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 349d062e74..9251b978af 100644
--- a/naga/tests/out/wgsl/buffer.frag.wgsl
+++ b/naga/tests/out/wgsl/buffer.frag.wgsl
@@ -20,7 +20,6 @@ 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 909ff07209..da6f754576 100644
--- a/naga/tests/out/wgsl/clamp-splat.vert.wgsl
+++ b/naga/tests/out/wgsl/clamp-splat.vert.wgsl
@@ -9,7 +9,6 @@ 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 477317594f..c72395367f 100644
--- a/naga/tests/out/wgsl/collatz.wgsl
+++ b/naga/tests/out/wgsl/collatz.wgsl
@@ -38,5 +38,4 @@ 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 20a43e2044..f8a53da0a4 100644
--- a/naga/tests/out/wgsl/const-exprs.wgsl
+++ b/naga/tests/out/wgsl/const-exprs.wgsl
@@ -44,7 +44,6 @@ fn non_constant_initializers() {
     let _e10 = y;
     let _e11 = z;
     out_3 = vec4<i32>(_e8, _e9, _e10, _e11);
-    return;
 }
 
 fn splat_of_constant() {
@@ -88,5 +87,4 @@ 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 dee49c2463..5b5bf6f348 100644
--- a/naga/tests/out/wgsl/const-global-swizzle.frag.wgsl
+++ b/naga/tests/out/wgsl/const-global-swizzle.frag.wgsl
@@ -14,7 +14,6 @@ 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 fe1c1c02f4..84aceb1e80 100644
--- a/naga/tests/out/wgsl/const_assert.wgsl
+++ b/naga/tests/out/wgsl/const_assert.wgsl
@@ -2,6 +2,5 @@ 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 46697cb7c2..56ebf9ceb7 100644
--- a/naga/tests/out/wgsl/constant-array-size.frag.wgsl
+++ b/naga/tests/out/wgsl/constant-array-size.frag.wgsl
@@ -32,7 +32,6 @@ fn function() -> vec4<f32> {
 }
 
 fn main_1() {
-    return;
 }
 
 @fragment 
diff --git a/naga/tests/out/wgsl/control-flow.wgsl b/naga/tests/out/wgsl/control-flow.wgsl
index ad071af58a..69ef57740d 100644
--- a/naga/tests/out/wgsl/control-flow.wgsl
+++ b/naga/tests/out/wgsl/control-flow.wgsl
@@ -27,7 +27,6 @@ fn loop_switch_continue(x: i32) {
             }
         }
     }
-    return;
 }
 
 fn loop_switch_continue_nesting(x_1: i32, y: i32, z: i32) {
@@ -74,7 +73,6 @@ 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) {
@@ -113,7 +111,6 @@ 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) 
@@ -160,18 +157,14 @@ 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/declarations.frag.wgsl b/naga/tests/out/wgsl/declarations.frag.wgsl
index c0c0236ad6..d8f47b0b00 100644
--- a/naga/tests/out/wgsl/declarations.frag.wgsl
+++ b/naga/tests/out/wgsl/declarations.frag.wgsl
@@ -51,7 +51,6 @@ 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 01c798586a..75eed3eec3 100644
--- a/naga/tests/out/wgsl/double-math-functions.frag.wgsl
+++ b/naga/tests/out/wgsl/double-math-functions.frag.wgsl
@@ -90,7 +90,6 @@ 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 9d1c6ac9d9..46b5369bef 100644
--- a/naga/tests/out/wgsl/empty.wgsl
+++ b/naga/tests/out/wgsl/empty.wgsl
@@ -1,4 +1,3 @@
 @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 ec53847d5f..f8632f103e 100644
--- a/naga/tests/out/wgsl/expressions.frag.wgsl
+++ b/naga/tests/out/wgsl/expressions.frag.wgsl
@@ -36,7 +36,6 @@ 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) {
@@ -58,7 +57,6 @@ 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) {
@@ -95,7 +93,6 @@ 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>) {
@@ -123,7 +120,6 @@ 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) {
@@ -160,7 +156,6 @@ 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>) {
@@ -188,7 +183,6 @@ 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>) {
@@ -217,7 +211,6 @@ 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>) {
@@ -257,7 +250,6 @@ 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>) {
@@ -276,7 +268,6 @@ fn testUnaryOpMat(a_16: mat3x3<f32>) {
     const _e12 = vec3(1f);
     a_17 = (_e10 - mat3x3<f32>(_e12, _e12, _e12));
     v_8 = _e10;
-    return;
 }
 
 fn testStructConstructor() {
@@ -295,7 +286,6 @@ fn testArrayConstructor() {
 }
 
 fn testFreestandingConstructor() {
-    return;
 }
 
 fn testNonImplicitCastVectorCast() {
@@ -304,11 +294,9 @@ 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) {
@@ -361,7 +349,6 @@ fn ternary(a_20: bool) {
     }
     let _e31 = local_4;
     nested = _e31;
-    return;
 }
 
 fn testMatrixMultiplication(a_22: mat4x3<f32>, b_18: mat4x4<f32>) {
@@ -374,14 +361,12 @@ 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>) {
@@ -400,7 +385,6 @@ fn indexConstantNonConstantIndex(i: i32) {
     let _e6 = i_1;
     let _e11 = local_5.array_[_e6];
     a_26 = _e11;
-    return;
 }
 
 fn testSwizzleWrites(a_27: vec3<f32>) {
@@ -419,7 +403,6 @@ 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() {
@@ -435,7 +418,6 @@ 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 65699237fe..c461908323 100644
--- a/naga/tests/out/wgsl/f64.wgsl
+++ b/naga/tests/out/wgsl/f64.wgsl
@@ -13,5 +13,4 @@ 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 57d11c1d24..69be9f8e1b 100644
--- a/naga/tests/out/wgsl/fma.frag.wgsl
+++ b/naga/tests/out/wgsl/fma.frag.wgsl
@@ -28,7 +28,6 @@ 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() {
@@ -37,7 +36,6 @@ 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 db7b81b146..0887cb513d 100644
--- a/naga/tests/out/wgsl/functions.wgsl
+++ b/naga/tests/out/wgsl/functions.wgsl
@@ -20,5 +20,4 @@ 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 62aaeb2033..c68f5fe0b1 100644
--- a/naga/tests/out/wgsl/functions_call.frag.wgsl
+++ b/naga/tests/out/wgsl/functions_call.frag.wgsl
@@ -1,5 +1,4 @@
 fn swizzleCallee(a: ptr<function, vec2<f32>>) {
-    return;
 }
 
 fn swizzleCaller(a_1: vec3<f32>) {
@@ -15,11 +14,9 @@ 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) {
@@ -30,11 +27,9 @@ 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>) {
@@ -49,11 +44,9 @@ 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/globals.wgsl b/naga/tests/out/wgsl/globals.wgsl
index 229d40ccf6..e85185c85f 100644
--- a/naga/tests/out/wgsl/globals.wgsl
+++ b/naga/tests/out/wgsl/globals.wgsl
@@ -23,7 +23,6 @@ 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_() {
@@ -67,5 +66,4 @@ 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 0c10eda857..21647f7fef 100644
--- a/naga/tests/out/wgsl/image.wgsl
+++ b/naga/tests/out/wgsl/image.wgsl
@@ -61,7 +61,6 @@ 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 6b57b52aeb..c91bb338e1 100644
--- a/naga/tests/out/wgsl/images.frag.wgsl
+++ b/naga/tests/out/wgsl/images.frag.wgsl
@@ -28,7 +28,6 @@ fn testImg1D(coord: i32) {
     let _e17 = coord_1;
     let _e18 = textureLoad(img1D, _e17);
     c = _e18;
-    return;
 }
 
 fn testImg1DArray(coord_2: vec2<i32>) {
@@ -45,7 +44,6 @@ 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>) {
@@ -61,7 +59,6 @@ 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>) {
@@ -78,7 +75,6 @@ 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>) {
@@ -94,7 +90,6 @@ 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>) {
@@ -108,7 +103,6 @@ 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>) {
@@ -120,7 +114,6 @@ 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>) {
@@ -130,11 +123,9 @@ 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 19d383eba5..900702812b 100644
--- a/naga/tests/out/wgsl/int64.wgsl
+++ b/naga/tests/out/wgsl/int64.wgsl
@@ -186,5 +186,4 @@ 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 1f46278ab8..d46197662c 100644
--- a/naga/tests/out/wgsl/interface.wgsl
+++ b/naga/tests/out/wgsl/interface.wgsl
@@ -35,7 +35,6 @@ 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 adb90568e4..c9d8a6f5c2 100644
--- a/naga/tests/out/wgsl/interpolate.wgsl
+++ b/naga/tests/out/wgsl/interpolate.wgsl
@@ -35,5 +35,4 @@ 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 aeaac63547..706f9ce7fb 100644
--- a/naga/tests/out/wgsl/interpolate_compat.wgsl
+++ b/naga/tests/out/wgsl/interpolate_compat.wgsl
@@ -33,5 +33,4 @@ 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 1efea1d9f6..325b7e7d40 100644
--- a/naga/tests/out/wgsl/inverse-polyfill.frag.wgsl
+++ b/naga/tests/out/wgsl/inverse-polyfill.frag.wgsl
@@ -25,7 +25,6 @@ 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 e787f96b10..fd9fa6aa92 100644
--- a/naga/tests/out/wgsl/lexical-scopes.wgsl
+++ b/naga/tests/out/wgsl/lexical-scopes.wgsl
@@ -1,23 +1,18 @@
 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) {
@@ -36,7 +31,6 @@ fn forLexicalScope(a_3: f32) {
             a_4 = (_e8 + 1i);
         }
     }
-    return;
 }
 
 fn whileLexicalScope(a_5: i32) {
@@ -48,7 +42,6 @@ fn whileLexicalScope(a_5: i32) {
         {
         }
     }
-    return;
 }
 
 fn switchLexicalScope(a_6: i32) {
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 a4bb55437a..26f56576ae 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,7 +19,6 @@ fn main_1() {
             i = (_e10 + 1i);
         }
     }
-    return;
 }
 
 @compute @workgroup_size(1, 1, 1) 
diff --git a/naga/tests/out/wgsl/math-functions.frag.wgsl b/naga/tests/out/wgsl/math-functions.frag.wgsl
index 10dc84c9a5..608a37a2bc 100644
--- a/naga/tests/out/wgsl/math-functions.frag.wgsl
+++ b/naga/tests/out/wgsl/math-functions.frag.wgsl
@@ -149,7 +149,6 @@ 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/module-scope.wgsl b/naga/tests/out/wgsl/module-scope.wgsl
index 48e3325e53..d9a41d3146 100644
--- a/naga/tests/out/wgsl/module-scope.wgsl
+++ b/naga/tests/out/wgsl/module-scope.wgsl
@@ -10,7 +10,6 @@ var Texture: texture_2d<f32>;
 var Sampler: sampler;
 
 fn statement() {
-    return;
 }
 
 fn returns() -> S {
diff --git a/naga/tests/out/wgsl/multiview.wgsl b/naga/tests/out/wgsl/multiview.wgsl
index 51192d2f7a..69dd53775c 100644
--- a/naga/tests/out/wgsl/multiview.wgsl
+++ b/naga/tests/out/wgsl/multiview.wgsl
@@ -1,4 +1,3 @@
 @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 2194a01df1..711b53c7a4 100644
--- a/naga/tests/out/wgsl/operators.wgsl
+++ b/naga/tests/out/wgsl/operators.wgsl
@@ -226,7 +226,6 @@ fn assignment() {
     vec0_[1i] = (_e37 + 1i);
     let _e41 = vec0_[1i];
     vec0_[1i] = (_e41 - 1i);
-    return;
 }
 
 fn negation_avoids_prefix_decrement() {
@@ -250,5 +249,4 @@ fn main(@builtin(workgroup_id) id: vec3<u32>) {
     bit();
     comparison();
     assignment();
-    return;
 }
diff --git a/naga/tests/out/wgsl/pointers.wgsl b/naga/tests/out/wgsl/pointers.wgsl
index 310f53cf08..42b0927f98 100644
--- a/naga/tests/out/wgsl/pointers.wgsl
+++ b/naga/tests/out/wgsl/pointers.wgsl
@@ -10,19 +10,16 @@ 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 15916303bb..6ce6114da8 100644
--- a/naga/tests/out/wgsl/prepostfix.frag.wgsl
+++ b/naga/tests/out/wgsl/prepostfix.frag.wgsl
@@ -29,7 +29,6 @@ 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 9de0a19c4b..23d82a6b93 100644
--- a/naga/tests/out/wgsl/quad_glsl.frag.wgsl
+++ b/naga/tests/out/wgsl/quad_glsl.frag.wgsl
@@ -7,7 +7,6 @@ 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 0a3d7cecac..ca7c705f9d 100644
--- a/naga/tests/out/wgsl/quad_glsl.vert.wgsl
+++ b/naga/tests/out/wgsl/quad_glsl.vert.wgsl
@@ -16,7 +16,6 @@ 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 0f79bd9693..332134287d 100644
--- a/naga/tests/out/wgsl/sampler-functions.frag.wgsl
+++ b/naga/tests/out/wgsl/sampler-functions.frag.wgsl
@@ -27,7 +27,6 @@ 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 170b0491b3..1dcece4ae2 100644
--- a/naga/tests/out/wgsl/samplers.frag.wgsl
+++ b/naga/tests/out/wgsl/samplers.frag.wgsl
@@ -118,7 +118,6 @@ fn testTex1D(coord: f32) {
     let _e179 = coord_1;
     let _e183 = textureLoad(tex1D, i32(_e179), 3i);
     c = _e183;
-    return;
 }
 
 fn testTex1DArray(coord_2: vec2<f32>) {
@@ -159,7 +158,6 @@ 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>) {
@@ -279,7 +277,6 @@ 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>) {
@@ -347,7 +344,6 @@ 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>) {
@@ -394,7 +390,6 @@ 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>) {
@@ -425,7 +420,6 @@ 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>) {
@@ -451,7 +445,6 @@ 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>) {
@@ -473,7 +466,6 @@ 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>) {
@@ -500,7 +492,6 @@ 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>) {
@@ -518,7 +509,6 @@ 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>) {
@@ -594,7 +584,6 @@ 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>) {
@@ -608,7 +597,6 @@ 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>) {
@@ -624,11 +612,9 @@ 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 bc36eb2075..8a25c1e1c2 100644
--- a/naga/tests/out/wgsl/statements.frag.wgsl
+++ b/naga/tests/out/wgsl/statements.frag.wgsl
@@ -54,7 +54,6 @@ 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 b1569bde3e..e03348c4d3 100644
--- a/naga/tests/out/wgsl/struct-layout.wgsl
+++ b/naga/tests/out/wgsl/struct-layout.wgsl
@@ -36,7 +36,6 @@ fn no_padding_comp() {
     x = _e2;
     let _e4 = no_padding_storage;
     x = _e4;
-    return;
 }
 
 @fragment 
@@ -57,5 +56,4 @@ 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 25f713b357..b1fa9c2511 100644
--- a/naga/tests/out/wgsl/subgroup-operations.wgsl
+++ b/naga/tests/out/wgsl/subgroup-operations.wgsl
@@ -27,5 +27,4 @@ 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/vector-functions.frag.wgsl b/naga/tests/out/wgsl/vector-functions.frag.wgsl
index 0f28a40c87..bb5e771411 100644
--- a/naga/tests/out/wgsl/vector-functions.frag.wgsl
+++ b/naga/tests/out/wgsl/vector-functions.frag.wgsl
@@ -28,7 +28,6 @@ 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>) {
@@ -61,7 +60,6 @@ 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>) {
@@ -94,7 +92,6 @@ 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>) {
@@ -127,7 +124,6 @@ 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>) {
@@ -153,11 +149,9 @@ 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 be85b033b9..10568d0e75 100644
--- a/naga/tests/out/wgsl/workgroup-uniform-load.wgsl
+++ b/naga/tests/out/wgsl/workgroup-uniform-load.wgsl
@@ -8,8 +8,5 @@ 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 fdad0477d6..0b41e4bb15 100644
--- a/naga/tests/out/wgsl/workgroup-var-init.wgsl
+++ b/naga/tests/out/wgsl/workgroup-var-init.wgsl
@@ -12,5 +12,4 @@ var<storage, read_write> output: array<u32, 512>;
 fn main() {
     let _e3 = w_mem.arr;
     output = _e3;
-    return;
 }
diff --git a/naga/tests/wgsl_errors.rs b/naga/tests/wgsl_errors.rs
index 968d6d59e0..a23eaf7047 100644
--- a/naga/tests/wgsl_errors.rs
+++ b/naga/tests/wgsl_errors.rs
@@ -2091,6 +2091,27 @@ fn function_param_redefinition_as_local() {
     )
 }
 
+#[test]
+fn function_must_return_value() {
+    check_validation!(
+        "fn func() -> i32 {
+        }":
+        Err(naga::valid::ValidationError::Function {
+            source: naga::valid::FunctionError::InvalidReturnType(_),
+            ..
+        })
+    );
+    check_validation!(
+        "fn func(x: i32) -> i32 {
+            let y = x + 10;
+        }":
+        Err(naga::valid::ValidationError::Function {
+            source: naga::valid::FunctionError::InvalidReturnType(_),
+            ..
+        })
+    );
+}
+
 #[test]
 fn constructor_type_error_span() {
     check(