diff --git a/CHANGELOG.md b/CHANGELOG.md index 7d9558c36b..3009a6af9b 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -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.2 (2025-07-23) ### Bug Fixes diff --git a/naga/src/back/spv/block.rs b/naga/src/back/spv/block.rs index 71487c6975..0cd414bfbe 100644 --- a/naga/src/back/spv/block.rs +++ b/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, diff --git a/naga/tests/in/wgsl/empty-if.toml b/naga/tests/in/wgsl/empty-if.toml new file mode 100644 index 0000000000..b7d70bcf51 --- /dev/null +++ b/naga/tests/in/wgsl/empty-if.toml @@ -0,0 +1,2 @@ +[spv] +version = [1, 6] diff --git a/naga/tests/in/wgsl/empty-if.wgsl b/naga/tests/in/wgsl/empty-if.wgsl new file mode 100644 index 0000000000..c72eb7a012 --- /dev/null +++ b/naga/tests/in/wgsl/empty-if.wgsl @@ -0,0 +1,9 @@ +@workgroup_size(1) +@compute +fn comp(@builtin(global_invocation_id) id: vec3) { + if (id.x == 0) { + + } + _ = 1+1; // otherwise, naga generates returns in the if statement. + return; +} \ No newline at end of file diff --git a/naga/tests/out/glsl/wgsl-empty-if.comp.Compute.glsl b/naga/tests/out/glsl/wgsl-empty-if.comp.Compute.glsl new file mode 100644 index 0000000000..9300c2802e --- /dev/null +++ b/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; +} + diff --git a/naga/tests/out/hlsl/wgsl-empty-if.hlsl b/naga/tests/out/hlsl/wgsl-empty-if.hlsl new file mode 100644 index 0000000000..6fd5db5e92 --- /dev/null +++ b/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; +} diff --git a/naga/tests/out/hlsl/wgsl-empty-if.ron b/naga/tests/out/hlsl/wgsl-empty-if.ron new file mode 100644 index 0000000000..03fe98f84c --- /dev/null +++ b/naga/tests/out/hlsl/wgsl-empty-if.ron @@ -0,0 +1,12 @@ +( + vertex:[ + ], + fragment:[ + ], + compute:[ + ( + entry_point:"comp", + target_profile:"cs_5_1", + ), + ], +) diff --git a/naga/tests/out/msl/wgsl-empty-if.msl b/naga/tests/out/msl/wgsl-empty-if.msl new file mode 100644 index 0000000000..d9ce803d61 --- /dev/null +++ b/naga/tests/out/msl/wgsl-empty-if.msl @@ -0,0 +1,16 @@ +// language: metal1.0 +#include +#include + +using metal::uint; + + +struct compInput { +}; +kernel void comp( + metal::uint3 id [[thread_position_in_grid]] +) { + if (id.x == 0u) { + } + return; +} diff --git a/naga/tests/out/spv/wgsl-access.spvasm b/naga/tests/out/spv/wgsl-access.spvasm index 78d263d5fc..31e8e5d4c0 100644 --- a/naga/tests/out/spv/wgsl-access.spvasm +++ b/naga/tests/out/spv/wgsl-access.spvasm @@ -1,16 +1,16 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 416 +; Bound: 414 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint Vertex %333 "foo_vert" %328 %331 -OpEntryPoint Fragment %389 "foo_frag" %388 -OpEntryPoint GLCompute %407 "foo_compute" -OpExecutionMode %389 OriginUpperLeft -OpExecutionMode %407 LocalSize 1 1 1 +OpEntryPoint Vertex %331 "foo_vert" %326 %329 +OpEntryPoint Fragment %387 "foo_frag" %386 +OpEntryPoint GLCompute %405 "foo_compute" +OpExecutionMode %387 OriginUpperLeft +OpExecutionMode %405 LocalSize 1 1 1 %3 = OpString "access.wgsl" OpSource Unknown 0 %3 "// This snapshot tests accessing various containers, dereferencing pointers. @@ -334,16 +334,16 @@ OpName %275 "a" OpName %284 "member_ptr" OpName %288 "s" OpName %294 "let_members_of_members" -OpName %306 "var_members_of_members" -OpName %307 "thing" -OpName %309 "inner" -OpName %312 "delishus" -OpName %328 "vi" -OpName %333 "foo_vert" -OpName %344 "foo" -OpName %345 "c2" -OpName %389 "foo_frag" -OpName %407 "foo_compute" +OpName %305 "var_members_of_members" +OpName %306 "thing" +OpName %308 "inner" +OpName %311 "delishus" +OpName %326 "vi" +OpName %331 "foo_vert" +OpName %342 "foo" +OpName %343 "c2" +OpName %387 "foo_frag" +OpName %405 "foo_compute" OpMemberDecorate %7 0 Offset 0 OpMemberDecorate %7 1 Offset 16 OpMemberDecorate %7 2 Offset 28 @@ -395,9 +395,9 @@ OpDecorate %62 DescriptorSet 0 OpDecorate %62 Binding 3 OpDecorate %63 Block OpMemberDecorate %63 0 Offset 0 -OpDecorate %328 BuiltIn VertexIndex -OpDecorate %331 BuiltIn Position -OpDecorate %388 Location 0 +OpDecorate %326 BuiltIn VertexIndex +OpDecorate %329 BuiltIn Position +OpDecorate %386 Location 0 %2 = OpTypeVoid %4 = OpTypeInt 32 0 %5 = OpTypeVector %4 3 @@ -533,45 +533,45 @@ OpDecorate %388 Location 0 %287 = OpConstantComposite %45 %286 %289 = OpTypePointer Function %45 %295 = OpConstantNull %47 -%308 = OpTypePointer Function %47 -%310 = OpTypePointer Function %46 -%311 = OpConstantNull %46 -%313 = OpConstantNull %6 -%329 = OpTypePointer Input %4 -%328 = OpVariable %329 Input -%332 = OpTypePointer Output %32 -%331 = OpVariable %332 Output -%335 = OpTypePointer StorageBuffer %24 -%338 = OpConstant %9 0 -%339 = OpConstant %4 3 -%340 = OpConstant %6 3 -%341 = OpConstant %6 4 -%342 = OpConstant %6 5 -%343 = OpConstantNull %30 -%346 = OpTypePointer Function %33 -%347 = OpConstantNull %33 -%353 = OpTypePointer StorageBuffer %10 -%356 = OpTypePointer StorageBuffer %19 -%359 = OpTypePointer StorageBuffer %11 -%360 = OpTypePointer StorageBuffer %9 -%363 = OpTypePointer StorageBuffer %20 -%366 = OpTypePointer StorageBuffer %8 -%367 = OpTypePointer StorageBuffer %6 -%372 = OpConstant %9 -2147483600 -%373 = OpConstant %9 2147483500 -%382 = OpTypeVector %6 4 -%388 = OpVariable %332 Output -%391 = OpConstantComposite %11 %338 %338 %338 -%392 = OpConstantComposite %11 %71 %71 %71 -%393 = OpConstantComposite %11 %73 %73 %73 -%394 = OpConstantComposite %11 %75 %75 %75 -%395 = OpConstantComposite %10 %391 %392 %393 %394 -%396 = OpConstantComposite %18 %48 %48 -%397 = OpConstantComposite %18 %44 %44 -%398 = OpConstantComposite %19 %396 %397 -%399 = OpConstantNull %24 -%400 = OpConstantComposite %32 %338 %338 %338 %338 -%408 = OpConstantTrue %42 +%307 = OpTypePointer Function %47 +%309 = OpTypePointer Function %46 +%310 = OpConstantNull %46 +%312 = OpConstantNull %6 +%327 = OpTypePointer Input %4 +%326 = OpVariable %327 Input +%330 = OpTypePointer Output %32 +%329 = OpVariable %330 Output +%333 = OpTypePointer StorageBuffer %24 +%336 = OpConstant %9 0 +%337 = OpConstant %4 3 +%338 = OpConstant %6 3 +%339 = OpConstant %6 4 +%340 = OpConstant %6 5 +%341 = OpConstantNull %30 +%344 = OpTypePointer Function %33 +%345 = OpConstantNull %33 +%351 = OpTypePointer StorageBuffer %10 +%354 = OpTypePointer StorageBuffer %19 +%357 = OpTypePointer StorageBuffer %11 +%358 = OpTypePointer StorageBuffer %9 +%361 = OpTypePointer StorageBuffer %20 +%364 = OpTypePointer StorageBuffer %8 +%365 = OpTypePointer StorageBuffer %6 +%370 = OpConstant %9 -2147483600 +%371 = OpConstant %9 2147483500 +%380 = OpTypeVector %6 4 +%386 = OpVariable %330 Output +%389 = OpConstantComposite %11 %336 %336 %336 +%390 = OpConstantComposite %11 %71 %71 %71 +%391 = OpConstantComposite %11 %73 %73 %73 +%392 = OpConstantComposite %11 %75 %75 %75 +%393 = OpConstantComposite %10 %389 %390 %391 %392 +%394 = OpConstantComposite %18 %48 %48 +%395 = OpConstantComposite %18 %44 %44 +%396 = OpConstantComposite %19 %394 %395 +%397 = OpConstantNull %24 +%398 = OpConstantComposite %32 %336 %336 %336 %336 +%406 = OpConstantTrue %42 %66 = OpFunction %2 None %67 %65 = OpLabel %94 = OpVariable %95 Function %70 @@ -939,164 +939,158 @@ OpLine %3 228 9 %300 = OpBitcast %4 %298 %301 = OpINotEqual %42 %299 %300 OpLine %3 228 5 -OpSelectionMerge %302 None -OpBranchConditional %301 %302 %302 -%302 = OpLabel OpLine %3 232 12 -%303 = OpCompositeExtract %46 %295 0 -%304 = OpCompositeExtract %6 %303 0 -OpReturnValue %304 +%302 = OpCompositeExtract %46 %295 0 +%303 = OpCompositeExtract %6 %302 0 +OpReturnValue %303 OpFunctionEnd -%306 = OpFunction %6 None %285 -%305 = OpLabel -%307 = OpVariable %308 Function %295 -%309 = OpVariable %310 Function %311 -%312 = OpVariable %95 Function %313 -OpBranch %314 -%314 = OpLabel +%305 = OpFunction %6 None %285 +%304 = OpLabel +%306 = OpVariable %307 Function %295 +%308 = OpVariable %309 Function %310 +%311 = OpVariable %95 Function %312 +OpBranch %313 +%313 = OpLabel OpLine %3 238 17 -%315 = OpAccessChain %310 %307 %48 -%316 = OpLoad %46 %315 +%314 = OpAccessChain %309 %306 %48 +%315 = OpLoad %46 %314 OpLine %3 238 5 -OpStore %309 %316 +OpStore %308 %315 OpLine %3 239 20 -%317 = OpAccessChain %95 %309 %48 -%318 = OpLoad %6 %317 +%316 = OpAccessChain %95 %308 %48 +%317 = OpLoad %6 %316 OpLine %3 239 5 -OpStore %312 %318 +OpStore %311 %317 OpLine %3 241 9 -%319 = OpAccessChain %34 %307 %44 -%320 = OpLoad %4 %319 -%321 = OpLoad %6 %312 -%322 = OpBitcast %4 %321 -%323 = OpINotEqual %42 %320 %322 +%318 = OpAccessChain %34 %306 %44 +%319 = OpLoad %4 %318 +%320 = OpLoad %6 %311 +%321 = OpBitcast %4 %320 +%322 = OpINotEqual %42 %319 %321 OpLine %3 241 5 -OpSelectionMerge %324 None -OpBranchConditional %323 %324 %324 -%324 = OpLabel OpLine %3 245 12 -%325 = OpAccessChain %95 %307 %48 %48 -%326 = OpLoad %6 %325 -OpReturnValue %326 +%323 = OpAccessChain %95 %306 %48 %48 +%324 = OpLoad %6 %323 +OpReturnValue %324 OpFunctionEnd -%333 = OpFunction %2 None %67 -%327 = OpLabel -%344 = OpVariable %28 Function %338 -%345 = OpVariable %346 Function %347 -%330 = OpLoad %4 %328 -%334 = OpAccessChain %68 %56 %48 -%336 = OpAccessChain %335 %59 %48 -%337 = OpAccessChain %141 %62 %48 -OpBranch %348 -%348 = OpLabel +%331 = OpFunction %2 None %67 +%325 = OpLabel +%342 = OpVariable %28 Function %336 +%343 = OpVariable %344 Function %345 +%328 = OpLoad %4 %326 +%332 = OpAccessChain %68 %56 %48 +%334 = OpAccessChain %333 %59 %48 +%335 = OpAccessChain %141 %62 %48 +OpBranch %346 +%346 = OpLabel OpLine %3 1 1 -%349 = OpLoad %9 %344 +%347 = OpLoad %9 %342 OpLine %3 118 5 -OpStore %344 %71 +OpStore %342 %71 OpLine %3 120 9 -%350 = OpLoad %7 %52 +%348 = OpLoad %7 %52 OpLine %3 121 5 -%351 = OpFunctionCall %2 %66 +%349 = OpFunctionCall %2 %66 OpLine %3 122 5 -%352 = OpFunctionCall %2 %140 +%350 = OpFunctionCall %2 %140 OpLine %3 125 19 -%354 = OpAccessChain %353 %54 %48 -%355 = OpLoad %10 %354 +%352 = OpAccessChain %351 %54 %48 +%353 = OpLoad %10 %352 OpLine %3 126 15 -%357 = OpAccessChain %356 %54 %40 -%358 = OpLoad %19 %357 +%355 = OpAccessChain %354 %54 %40 +%356 = OpLoad %19 %355 OpLine %3 128 13 -%361 = OpAccessChain %360 %54 %48 %339 %48 -%362 = OpLoad %9 %361 +%359 = OpAccessChain %358 %54 %48 %337 %48 +%360 = OpLoad %9 %359 OpLine %3 129 13 OpLine %3 129 22 -%364 = OpArrayLength %4 %54 5 +%362 = OpArrayLength %4 %54 5 OpLine %3 129 13 -%365 = OpISub %4 %364 %15 -%368 = OpAccessChain %367 %54 %31 %365 %48 -%369 = OpLoad %6 %368 +%363 = OpISub %4 %362 %15 +%366 = OpAccessChain %365 %54 %31 %363 %48 +%367 = OpLoad %6 %366 OpLine %3 130 13 -%370 = OpLoad %24 %336 +%368 = OpLoad %24 %334 OpLine %3 133 56 OpLine %3 133 56 OpLine %3 134 21 -%371 = OpFunctionCall %9 %198 %344 +%369 = OpFunctionCall %9 %198 %342 OpLine %3 137 31 -%374 = OpExtInst %9 %1 FClamp %362 %372 %373 -%375 = OpConvertFToS %6 %374 +%372 = OpExtInst %9 %1 FClamp %360 %370 %371 +%373 = OpConvertFToS %6 %372 OpLine %3 137 14 -%376 = OpCompositeConstruct %33 %369 %375 %340 %341 %342 +%374 = OpCompositeConstruct %33 %367 %373 %338 %339 %340 OpLine %3 137 5 -OpStore %345 %376 +OpStore %343 %374 OpLine %3 138 5 -%377 = OpIAdd %4 %330 %44 +%375 = OpIAdd %4 %328 %44 OpLine %3 138 5 -%378 = OpAccessChain %95 %345 %377 -OpStore %378 %286 +%376 = OpAccessChain %95 %343 %375 +OpStore %376 %286 OpLine %3 139 17 -%379 = OpAccessChain %95 %345 %330 -%380 = OpLoad %6 %379 +%377 = OpAccessChain %95 %343 %328 +%378 = OpLoad %6 %377 OpLine %3 141 5 -%381 = OpFunctionCall %9 %204 %343 +%379 = OpFunctionCall %9 %204 %341 OpLine %3 143 22 -%383 = OpCompositeConstruct %382 %380 %380 %380 %380 -%384 = OpConvertSToF %32 %383 -%385 = OpMatrixTimesVector %11 %355 %384 +%381 = OpCompositeConstruct %380 %378 %378 %378 %378 +%382 = OpConvertSToF %32 %381 +%383 = OpMatrixTimesVector %11 %353 %382 OpLine %3 143 12 -%386 = OpCompositeConstruct %32 %385 %73 -OpStore %331 %386 +%384 = OpCompositeConstruct %32 %383 %73 +OpStore %329 %384 OpReturn OpFunctionEnd -%389 = OpFunction %2 None %67 -%387 = OpLabel -%390 = OpAccessChain %335 %59 %48 -OpBranch %401 -%401 = OpLabel +%387 = OpFunction %2 None %67 +%385 = OpLabel +%388 = OpAccessChain %333 %59 %48 +OpBranch %399 +%399 = OpLabel OpLine %3 149 5 OpLine %3 149 5 OpLine %3 149 5 -%402 = OpAccessChain %360 %54 %48 %44 %15 -OpStore %402 %71 +%400 = OpAccessChain %358 %54 %48 %44 %15 +OpStore %400 %71 OpLine %3 150 5 OpLine %3 150 31 OpLine %3 150 47 OpLine %3 150 63 OpLine %3 150 19 OpLine %3 150 5 -%403 = OpAccessChain %353 %54 %48 -OpStore %403 %395 +%401 = OpAccessChain %351 %54 %48 +OpStore %401 %393 OpLine %3 151 5 OpLine %3 151 35 OpLine %3 151 15 OpLine %3 151 5 -%404 = OpAccessChain %356 %54 %40 -OpStore %404 %398 +%402 = OpAccessChain %354 %54 %40 +OpStore %402 %396 OpLine %3 152 5 OpLine %3 152 5 OpLine %3 152 5 -%405 = OpAccessChain %367 %54 %31 %44 %48 -OpStore %405 %70 +%403 = OpAccessChain %365 %54 %31 %44 %48 +OpStore %403 %70 OpLine %3 153 5 -OpStore %390 %399 +OpStore %388 %397 OpLine %3 155 12 -OpStore %388 %400 +OpStore %386 %398 OpReturn OpFunctionEnd -%407 = OpFunction %2 None %67 -%406 = OpLabel -OpBranch %409 -%409 = OpLabel +%405 = OpFunction %2 None %67 +%404 = OpLabel +OpBranch %407 +%407 = OpLabel OpLine %3 250 5 -%410 = OpFunctionCall %2 %224 +%408 = OpFunctionCall %2 %224 OpLine %3 251 5 -%411 = OpFunctionCall %2 %261 +%409 = OpFunctionCall %2 %261 OpLine %3 252 5 -%412 = OpFunctionCall %42 %273 %408 +%410 = OpFunctionCall %42 %273 %406 OpLine %3 253 5 -%413 = OpFunctionCall %6 %284 +%411 = OpFunctionCall %6 %284 OpLine %3 254 5 -%414 = OpFunctionCall %6 %294 +%412 = OpFunctionCall %6 %294 OpLine %3 255 5 -%415 = OpFunctionCall %6 %306 +%413 = OpFunctionCall %6 %305 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/spv/wgsl-empty-if.spvasm b/naga/tests/out/spv/wgsl-empty-if.spvasm new file mode 100644 index 0000000000..3845afcede --- /dev/null +++ b/naga/tests/out/spv/wgsl-empty-if.spvasm @@ -0,0 +1,29 @@ +; SPIR-V +; Version: 1.6 +; Generator: rspirv +; Bound: 18 +OpCapability Shader +%1 = OpExtInstImport "GLSL.std.450" +OpMemoryModel Logical GLSL450 +OpEntryPoint GLCompute %9 "comp" %6 +OpExecutionMode %9 LocalSize 1 1 1 +OpDecorate %6 BuiltIn GlobalInvocationId +%2 = OpTypeVoid +%4 = OpTypeInt 32 0 +%3 = OpTypeVector %4 3 +%7 = OpTypePointer Input %3 +%6 = OpVariable %7 Input +%10 = OpTypeFunction %2 +%11 = OpConstant %4 0 +%12 = OpTypeInt 32 1 +%13 = OpConstant %12 2 +%16 = OpTypeBool +%9 = OpFunction %2 None %10 +%5 = OpLabel +%8 = OpLoad %3 %6 +OpBranch %14 +%14 = OpLabel +%15 = OpCompositeExtract %4 %8 0 +%17 = OpIEqual %16 %15 %11 +OpReturn +OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/wgsl/wgsl-empty-if.wgsl b/naga/tests/out/wgsl/wgsl-empty-if.wgsl new file mode 100644 index 0000000000..5435020695 --- /dev/null +++ b/naga/tests/out/wgsl/wgsl-empty-if.wgsl @@ -0,0 +1,6 @@ +@compute @workgroup_size(1, 1, 1) +fn comp(@builtin(global_invocation_id) id: vec3) { + if (id.x == 0u) { + } + return; +}