Skip to content

Fix empty ifs causing errors on spirv 1.6 #7883

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 3 commits into
base: trunk
Choose a base branch
from
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 6 additions & 0 deletions CHANGELOG.md
Original file line number Diff line number Diff line change
Expand Up @@ -59,6 +59,12 @@ By @Vecvec in [#7913](https://github.com/gfx-rs/wgpu/pull/7913).

Naga now requires that no type be larger than 1 GB. This limit may be lowered in the future; feedback on an appropriate value for the limit is welcome. By @andyleiserson in [#7950](https://github.com/gfx-rs/wgpu/pull/7950).

### 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.1 (2025-07-10)

### Bug Fixes
Expand Down
109 changes: 58 additions & 51 deletions naga/src/back/spv/block.rs
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Expand Down
2 changes: 2 additions & 0 deletions naga/tests/in/wgsl/empty-if.toml
Original file line number Diff line number Diff line change
@@ -0,0 +1,2 @@
[spv]
version = [1, 6]
9 changes: 9 additions & 0 deletions naga/tests/in/wgsl/empty-if.wgsl
Original file line number Diff line number Diff line change
@@ -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;
}
15 changes: 15 additions & 0 deletions naga/tests/out/glsl/wgsl-empty-if.comp.Compute.glsl
Original file line number Diff line number Diff line change
@@ -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;
}

7 changes: 7 additions & 0 deletions naga/tests/out/hlsl/wgsl-empty-if.hlsl
Original file line number Diff line number Diff line change
@@ -0,0 +1,7 @@
[numthreads(1, 1, 1)]
void comp(uint3 id : SV_DispatchThreadID)
{
if ((id.x == 0u)) {
}
return;
}
12 changes: 12 additions & 0 deletions naga/tests/out/hlsl/wgsl-empty-if.ron
Original file line number Diff line number Diff line change
@@ -0,0 +1,12 @@
(
vertex:[
],
fragment:[
],
compute:[
(
entry_point:"comp",
target_profile:"cs_5_1",
),
],
)
16 changes: 16 additions & 0 deletions naga/tests/out/msl/wgsl-empty-if.msl
Original file line number Diff line number Diff line change
@@ -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;
}
Loading