Fix empty ifs causing errors on spirv 1.6 (#7883)

Author: Vecvec

CHANGELOG.md

@@ -109,6 +109,12 @@ By @Vecvec in [#7913](https://github.com/gfx-rs/wgpu/pull/7913).
 
 - Fixed memory leak in vulkan backend. By @cwfitzgerald in [#8031](https://github.com/gfx-rs/wgpu/pull/8031).
 
+### Bug Fixes
+
+#### Naga
+
+- Fix empty `if` statements causing errors on spirv 1.6+. By @Vecvec in [#7883](https://github.com/gfx-rs/wgpu/pull/7883).
+
 ## v26.0.2 (2025-07-23)
 
 ### Bug Fixes

naga/src/back/spv/block.rs

@@ -2979,62 +2979,69 @@ impl BlockContext<'_> {
                     ref accept,
                     ref reject,
                 } => {
-                    let condition_id = self.cached[condition];
+                    // In spirv 1.6, in a conditional branch the two block ids
+                    // of the branches can't have the same label. If `accept`
+                    // and `reject` are both empty (e.g. in `if (condition) {}`)
+                    // merge id will be both labels. Because both branches are
+                    // empty, we can skip the if statement.
+                    if !(accept.is_empty() && reject.is_empty()) {
+                        let condition_id = self.cached[condition];
+
+                        let merge_id = self.gen_id();
+                        block.body.push(Instruction::selection_merge(
+                            merge_id,
+                            spirv::SelectionControl::NONE,
+                        ));
 
-                    let merge_id = self.gen_id();
-                    block.body.push(Instruction::selection_merge(
-                        merge_id,
-                        spirv::SelectionControl::NONE,
-                    ));
+                        let accept_id = if accept.is_empty() {
+                            None
+                        } else {
+                            Some(self.gen_id())
+                        };
+                        let reject_id = if reject.is_empty() {
+                            None
+                        } else {
+                            Some(self.gen_id())
+                        };
 
-                    let accept_id = if accept.is_empty() {
-                        None
-                    } else {
-                        Some(self.gen_id())
-                    };
-                    let reject_id = if reject.is_empty() {
-                        None
-                    } else {
-                        Some(self.gen_id())
-                    };
+                        self.function.consume(
+                            block,
+                            Instruction::branch_conditional(
+                                condition_id,
+                                accept_id.unwrap_or(merge_id),
+                                reject_id.unwrap_or(merge_id),
+                            ),
+                        );
 
-                    self.function.consume(
-                        block,
-                        Instruction::branch_conditional(
-                            condition_id,
-                            accept_id.unwrap_or(merge_id),
-                            reject_id.unwrap_or(merge_id),
-                        ),
-                    );
+                        if let Some(block_id) = accept_id {
+                            // We can ignore the `BlockExitDisposition` returned here because,
+                            // even if `merge_id` is not actually reachable, it is always
+                            // referred to by the `OpSelectionMerge` instruction we emitted
+                            // earlier.
+                            let _ = self.write_block(
+                                block_id,
+                                accept,
+                                BlockExit::Branch { target: merge_id },
+                                loop_context,
+                                debug_info,
+                            )?;
+                        }
+                        if let Some(block_id) = reject_id {
+                            // We can ignore the `BlockExitDisposition` returned here because,
+                            // even if `merge_id` is not actually reachable, it is always
+                            // referred to by the `OpSelectionMerge` instruction we emitted
+                            // earlier.
+                            let _ = self.write_block(
+                                block_id,
+                                reject,
+                                BlockExit::Branch { target: merge_id },
+                                loop_context,
+                                debug_info,
+                            )?;
+                        }
 
-                    if let Some(block_id) = accept_id {
-                        // We can ignore the `BlockExitDisposition` returned here because,
-                        // even if `merge_id` is not actually reachable, it is always
-                        // referred to by the `OpSelectionMerge` instruction we emitted
-                        // earlier.
-                        let _ = self.write_block(
-                            block_id,
-                            accept,
-                            BlockExit::Branch { target: merge_id },
-                            loop_context,
-                            debug_info,
-                        )?;
-                    }
-                    if let Some(block_id) = reject_id {
-                        // We can ignore the `BlockExitDisposition` returned here because,
-                        // even if `merge_id` is not actually reachable, it is always
-                        // referred to by the `OpSelectionMerge` instruction we emitted
-                        // earlier.
-                        let _ = self.write_block(
-                            block_id,
-                            reject,
-                            BlockExit::Branch { target: merge_id },
-                            loop_context,
-                            debug_info,
-                        )?;
+                        block = Block::new(merge_id);
                     }
-
-                    block = Block::new(merge_id);
                 }
                 Statement::Switch {
                     selector,

naga/tests/in/wgsl/empty-if.toml

@@ -0,0 +1,2 @@
+[spv]
+version = [1, 6]

naga/tests/in/wgsl/empty-if.wgsl

@@ -0,0 +1,9 @@
+@workgroup_size(1)
+@compute
+fn comp(@builtin(global_invocation_id) id: vec3<u32>) {
+    if (id.x == 0) {
+
+    }
+    _ = 1+1; // otherwise, naga generates returns in the if statement.
+    return;
+}

naga/tests/out/glsl/wgsl-empty-if.comp.Compute.glsl

@@ -0,0 +1,15 @@
+#version 310 es
+
+precision highp float;
+precision highp int;
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+
+void main() {
+    uvec3 id = gl_GlobalInvocationID;
+    if ((id.x == 0u)) {
+    }
+    return;
+}
+

naga/tests/out/hlsl/wgsl-empty-if.hlsl

@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void comp(uint3 id : SV_DispatchThreadID)
+{
+    if ((id.x == 0u)) {
+    }
+    return;
+}

naga/tests/out/hlsl/wgsl-empty-if.ron

@@ -0,0 +1,12 @@
+(
+    vertex:[
+    ],
+    fragment:[
+    ],
+    compute:[
+        (
+            entry_point:"comp",
+            target_profile:"cs_5_1",
+        ),
+    ],
+)

naga/tests/out/msl/wgsl-empty-if.msl

@@ -0,0 +1,16 @@
+// language: metal1.0
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using metal::uint;
+
+
+struct compInput {
+};
+kernel void comp(
+  metal::uint3 id [[thread_position_in_grid]]
+) {
+    if (id.x == 0u) {
+    }
+    return;
+}