From cb0ad2504ce8bbd868237032f49c10c944b490cb Mon Sep 17 00:00:00 2001 From: Jim Blandy Date: Fri, 27 Aug 2021 15:44:57 -0700 Subject: [PATCH] [spv-out] Handle break and continue in switch statements. Fixes #1030. Fixes #1017. --- src/back/spv/block.rs | 9 +- tests/in/control-flow.wgsl | 27 ++++ tests/out/glsl/control-flow.main.Compute.glsl | 25 ++++ tests/out/hlsl/control-flow.hlsl | 33 +++++ tests/out/msl/control-flow.msl | 39 ++++++ tests/out/spv/control-flow.spvasm | 120 +++++++++++++++--- tests/out/wgsl/control-flow.wgsl | 28 ++++ tests/snapshots.rs | 4 +- 8 files changed, 260 insertions(+), 25 deletions(-) diff --git a/src/back/spv/block.rs b/src/back/spv/block.rs index 3009085cee..e9970c23ca 100644 --- a/src/back/spv/block.rs +++ b/src/back/spv/block.rs @@ -1013,6 +1013,11 @@ impl<'w> BlockContext<'w> { Instruction::switch(selector_id, default_id, &raw_cases), ); + let inner_context = LoopContext { + break_id: Some(merge_id), + ..loop_context + }; + for (i, (case, raw_case)) in cases.iter().zip(raw_cases.iter()).enumerate() { let case_finish_id = if case.fall_through { match raw_cases.get(i + 1) { @@ -1026,11 +1031,11 @@ impl<'w> BlockContext<'w> { raw_case.label_id, &case.body, Some(case_finish_id), - LoopContext::default(), + inner_context, )?; } - self.write_block(default_id, default, Some(merge_id), LoopContext::default())?; + self.write_block(default_id, default, Some(merge_id), inner_context)?; block = Block::new(merge_id); } diff --git a/tests/in/control-flow.wgsl b/tests/in/control-flow.wgsl index 6cb0f3619b..50e8b22291 100644 --- a/tests/in/control-flow.wgsl +++ b/tests/in/control-flow.wgsl @@ -30,3 +30,30 @@ fn main([[builtin(global_invocation_id)]] global_id: vec3) { } } } + +fn switch_default_break(i: i32) { + switch (i) { + default: { + break; + } + } +} + +fn switch_case_break() { + switch(0) { + case 0: { + break; + } + } + return; +} + +fn loop_switch_continue(x: i32) { + loop { + switch (x) { + case 1: { + continue; + } + } + } +} diff --git a/tests/out/glsl/control-flow.main.Compute.glsl b/tests/out/glsl/control-flow.main.Compute.glsl index fd12e43461..2a42aa55c2 100644 --- a/tests/out/glsl/control-flow.main.Compute.glsl +++ b/tests/out/glsl/control-flow.main.Compute.glsl @@ -6,6 +6,31 @@ precision highp int; layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void switch_default_break(int i) { + switch(i) { + default: + break; + } +} + +void switch_case_break() { + switch(0) { + case 0: + break; + } + return; +} + +void loop_switch_continue(int x) { + while(true) { + switch(x) { + case 1: + continue; + } + } + return; +} + void main() { uvec3 global_id = gl_GlobalInvocationID; int pos = 0; diff --git a/tests/out/hlsl/control-flow.hlsl b/tests/out/hlsl/control-flow.hlsl index 020e228f0f..3dd0ab7a24 100644 --- a/tests/out/hlsl/control-flow.hlsl +++ b/tests/out/hlsl/control-flow.hlsl @@ -1,4 +1,37 @@ +void switch_default_break(int i) +{ + switch(i) { + default: { + break; + } + } +} + +void switch_case_break() +{ + switch(0) { + case 0: { + break; + break; + } + } + return; +} + +void loop_switch_continue(int x) +{ + while(true) { + switch(x) { + case 1: { + continue; + break; + } + } + } + return; +} + [numthreads(1, 1, 1)] void main(uint3 global_id : SV_DispatchThreadID) { diff --git a/tests/out/msl/control-flow.msl b/tests/out/msl/control-flow.msl index bd9b8d27d9..a2a80ac7bb 100644 --- a/tests/out/msl/control-flow.msl +++ b/tests/out/msl/control-flow.msl @@ -3,6 +3,45 @@ #include +void switch_default_break( + int i +) { + switch(i) { + default: { + break; + } + } +} + +void switch_case_break( +) { + switch(0) { + case 0: { + break; + break; + } + default: { + } + } + return; +} + +void loop_switch_continue( + int x +) { + while(true) { + switch(x) { + case 1: { + continue; + break; + } + default: { + } + } + } + return; +} + struct main1Input { }; kernel void main1( diff --git a/tests/out/spv/control-flow.spvasm b/tests/out/spv/control-flow.spvasm index b46424b801..fc79a6c7b8 100644 --- a/tests/out/spv/control-flow.spvasm +++ b/tests/out/spv/control-flow.spvasm @@ -1,29 +1,109 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 16 +; Bound: 56 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %9 "main" %6 -OpExecutionMode %9 LocalSize 1 1 1 -OpDecorate %6 BuiltIn GlobalInvocationId +OpEntryPoint GLCompute %41 "main" %38 +OpExecutionMode %41 LocalSize 1 1 1 +OpDecorate %38 BuiltIn GlobalInvocationId %2 = OpTypeVoid -%4 = OpTypeInt 32 0 -%3 = OpTypeVector %4 3 -%7 = OpTypePointer Input %3 -%6 = OpVariable %7 Input -%10 = OpTypeFunction %2 -%12 = OpConstant %4 2 -%13 = OpConstant %4 1 -%14 = OpConstant %4 72 -%15 = OpConstant %4 264 -%9 = OpFunction %2 None %10 -%5 = OpLabel -%8 = OpLoad %3 %6 -OpBranch %11 -%11 = OpLabel -OpControlBarrier %12 %13 %14 -OpControlBarrier %12 %12 %15 +%4 = OpTypeInt 32 1 +%3 = OpConstant %4 1 +%5 = OpConstant %4 0 +%6 = OpConstant %4 2 +%7 = OpConstant %4 3 +%9 = OpTypeInt 32 0 +%8 = OpTypeVector %9 3 +%13 = OpTypeFunction %2 %4 +%19 = OpTypeFunction %2 +%36 = OpTypePointer Function %4 +%39 = OpTypePointer Input %8 +%38 = OpVariable %39 Input +%43 = OpConstant %9 2 +%44 = OpConstant %9 1 +%45 = OpConstant %9 72 +%46 = OpConstant %9 264 +%12 = OpFunction %2 None %13 +%11 = OpFunctionParameter %4 +%10 = OpLabel +OpBranch %14 +%14 = OpLabel +OpSelectionMerge %15 None +OpSwitch %11 %16 +%16 = OpLabel +OpBranch %15 +%15 = OpLabel +OpReturn +OpFunctionEnd +%18 = OpFunction %2 None %19 +%17 = OpLabel +OpBranch %20 +%20 = OpLabel +OpSelectionMerge %21 None +OpSwitch %5 %22 0 %23 +%23 = OpLabel +OpBranch %21 +%22 = OpLabel +OpBranch %21 +%21 = OpLabel +OpReturn +OpFunctionEnd +%26 = OpFunction %2 None %13 +%25 = OpFunctionParameter %4 +%24 = OpLabel +OpBranch %27 +%27 = OpLabel +OpBranch %28 +%28 = OpLabel +OpLoopMerge %29 %31 None +OpBranch %30 +%30 = OpLabel +OpSelectionMerge %32 None +OpSwitch %25 %33 1 %34 +%34 = OpLabel +OpBranch %31 +%33 = OpLabel +OpBranch %32 +%32 = OpLabel +OpBranch %31 +%31 = OpLabel +OpBranch %28 +%29 = OpLabel +OpReturn +OpFunctionEnd +%41 = OpFunction %2 None %19 +%37 = OpLabel +%35 = OpVariable %36 Function +%40 = OpLoad %8 %38 +OpBranch %42 +%42 = OpLabel +OpControlBarrier %43 %44 %45 +OpControlBarrier %43 %43 %46 +OpSelectionMerge %47 None +OpSwitch %3 %48 +%48 = OpLabel +OpStore %35 %3 +OpBranch %47 +%47 = OpLabel +%49 = OpLoad %4 %35 +OpSelectionMerge %50 None +OpSwitch %49 %51 1 %52 2 %53 3 %54 4 %55 +%52 = OpLabel +OpStore %35 %5 +OpBranch %50 +%53 = OpLabel +OpStore %35 %3 +OpReturn +%54 = OpLabel +OpStore %35 %6 +OpBranch %55 +%55 = OpLabel +OpReturn +%51 = OpLabel +OpStore %35 %7 +OpReturn +%50 = OpLabel OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/control-flow.wgsl b/tests/out/wgsl/control-flow.wgsl index 3eb55cfd80..903ad99236 100644 --- a/tests/out/wgsl/control-flow.wgsl +++ b/tests/out/wgsl/control-flow.wgsl @@ -1,3 +1,31 @@ +fn switch_default_break(i: i32) { + switch(i) { + default: { + break; + } + } +} + +fn switch_case_break() { + switch(0) { + case 0: { + break; + } + } + return; +} + +fn loop_switch_continue(x: i32) { + loop { + switch(x) { + case 1: { + continue; + } + } + } + return; +} + [[stage(compute), workgroup_size(1, 1, 1)]] fn main([[builtin(global_invocation_id)]] global_id: vec3) { var pos: i32; diff --git a/tests/snapshots.rs b/tests/snapshots.rs index b346cc77ff..d25aa0df65 100644 --- a/tests/snapshots.rs +++ b/tests/snapshots.rs @@ -448,9 +448,7 @@ fn convert_wgsl() { ), ( "control-flow", - // TODO: SPIRV https://github.com/gfx-rs/naga/issues/1017 - //Targets::SPIRV | - Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL, + Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL, ), ( "standard",