[spv-out] Handle break and continue in switch statements.

Fixes #1030. Fixes #1017.
gfx-rs · Aug 27, 2021 · cb0ad25 · cb0ad25
1 parent 28547e3
commit cb0ad25
Show file tree

Hide file tree

Showing 8 changed files with 260 additions and 25 deletions.
diff --git 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
@@ -30,3 +30,30 @@ fn main([[builtin(global_invocation_id)]] global_id: vec3<u32>) {
 		}
     }
 }
+
+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
@@ -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
@@ -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
@@ -3,6 +3,45 @@
 #include <simd/simd.h>
 
 
+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
@@ -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
diff --git 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<u32>) {
     var pos: i32;

diff --git 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",