| // Copyright 2025 The Dawn & Tint Authors |
| // |
| // Redistribution and use in source and binary forms, with or without |
| // modification, are permitted provided that the following conditions are met: |
| // |
| // 1. Redistributions of source code must retain the above copyright notice, this |
| // list of conditions and the following disclaimer. |
| // |
| // 2. Redistributions in binary form must reproduce the above copyright notice, |
| // this list of conditions and the following disclaimer in the documentation |
| // and/or other materials provided with the distribution. |
| // |
| // 3. Neither the name of the copyright holder nor the names of its |
| // contributors may be used to endorse or promote products derived from |
| // this software without specific prior written permission. |
| // |
| // THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" |
| // AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE |
| // IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE |
| // DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE |
| // FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL |
| // DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR |
| // SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER |
| // CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, |
| // OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE |
| // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. |
| |
| #include "src/tint/lang/spirv/reader/parser/helper_test.h" |
| |
| namespace tint::spirv::reader { |
| namespace { |
| |
| TEST_F(SpirvParserTest, Branch_Forward) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %main_start = OpLabel |
| %1 = OpCopyObject %i32 %one |
| OpBranch %bb_2 |
| %bb_2 = OpLabel |
| %2 = OpCopyObject %i32 %two |
| OpBranch %bb_3 |
| %bb_3 = OpLabel |
| %3 = OpCopyObject %i32 %three |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = let 1i |
| %3:i32 = let 2i |
| %4:i32 = let 3i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %1 = OpCopyObject %i32 %one |
| OpSelectionMerge %bb_merge None |
| OpBranchConditional %true %bb_true %bb_false |
| %bb_true = OpLabel |
| %2 = OpCopyObject %i32 %two |
| OpReturn |
| %bb_false = OpLabel |
| %3 = OpCopyObject %i32 %three |
| OpBranch %bb_merge |
| %bb_merge = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = let 1i |
| if true [t: $B2, f: $B3] { # if_1 |
| $B2: { # true |
| %3:i32 = let 2i |
| ret |
| } |
| $B3: { # false |
| %4:i32 = let 3i |
| exit_if # if_1 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_Empty) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpSelectionMerge %99 None |
| OpBranchConditional %true %99 %99 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:bool = or true, true |
| if %2 [t: $B2, f: $B3] { # if_1 |
| $B2: { # true |
| exit_if # if_1 |
| } |
| $B3: { # false |
| unreachable |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_TrueToMerge) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %main_start = OpLabel |
| %1 = OpCopyObject %i32 %one |
| OpSelectionMerge %bb_merge None |
| OpBranchConditional %true %bb_merge %bb_false |
| %bb_false = OpLabel |
| %3 = OpCopyObject %i32 %three |
| OpBranch %bb_merge |
| %bb_merge = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = let 1i |
| if true [t: $B2, f: $B3] { # if_1 |
| $B2: { # true |
| exit_if # if_1 |
| } |
| $B3: { # false |
| %3:i32 = let 3i |
| exit_if # if_1 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_FalseToMerge) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %main_start = OpLabel |
| %1 = OpCopyObject %i32 %one |
| OpSelectionMerge %bb_merge None |
| OpBranchConditional %true %bb_true %bb_merge |
| %bb_true = OpLabel |
| %2 = OpCopyObject %i32 %two |
| OpReturn |
| %bb_merge = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = let 1i |
| if true [t: $B2, f: $B3] { # if_1 |
| $B2: { # true |
| %3:i32 = let 2i |
| ret |
| } |
| $B3: { # false |
| exit_if # if_1 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_TrueMatchesFalse) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %main_start = OpLabel |
| %1 = OpCopyObject %i32 %one |
| OpBranchConditional %true %bb_true %bb_true |
| %bb_true = OpLabel |
| %2 = OpCopyObject %i32 %two |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = let 1i |
| %3:bool = or true, true |
| if %3 [t: $B2, f: $B3] { # if_1 |
| $B2: { # true |
| %4:i32 = let 2i |
| ret |
| } |
| $B3: { # false |
| unreachable |
| } |
| } |
| unreachable |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_Nested_FalseExit) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %main_start = OpLabel |
| %1 = OpCopyObject %i32 %one |
| OpSelectionMerge %bb_merge None |
| OpBranchConditional %true %bb_true %bb_merge |
| %bb_true = OpLabel |
| OpBranchConditional %false %99 %bb_merge |
| %99 = OpLabel |
| %2 = OpCopyObject %i32 %two |
| OpBranch %bb_merge |
| %bb_merge = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = let 1i |
| if true [t: $B2, f: $B3] { # if_1 |
| $B2: { # true |
| if false [t: $B4, f: $B5] { # if_2 |
| $B4: { # true |
| %3:i32 = let 2i |
| exit_if # if_2 |
| } |
| $B5: { # false |
| exit_if # if_2 |
| } |
| } |
| exit_if # if_1 |
| } |
| $B3: { # false |
| exit_if # if_1 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_Nested_TrueExit) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %main_start = OpLabel |
| %1 = OpCopyObject %i32 %one |
| OpSelectionMerge %bb_merge None |
| OpBranchConditional %true %bb_true %bb_merge |
| %bb_true = OpLabel |
| OpBranchConditional %false %bb_merge %99 |
| %99 = OpLabel |
| %2 = OpCopyObject %i32 %two |
| OpBranch %bb_merge |
| %bb_merge = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = let 1i |
| if true [t: $B2, f: $B3] { # if_1 |
| $B2: { # true |
| if false [t: $B4, f: $B5] { # if_2 |
| $B4: { # true |
| exit_if # if_2 |
| } |
| $B5: { # false |
| %3:i32 = let 2i |
| exit_if # if_2 |
| } |
| } |
| exit_if # if_1 |
| } |
| $B3: { # false |
| exit_if # if_1 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_TrueBranchesToFalse) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %20 = OpLabel |
| OpSelectionMerge %99 None |
| OpBranchConditional %false %40 %30 |
| %30 = OpLabel |
| %50 = OpIAdd %i32 %one %one |
| OpReturn |
| %40 = OpLabel |
| %51 = OpIAdd %i32 %two %two |
| OpBranch %30 ; backtrack, but does dominate %30 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| if false [t: $B2, f: $B3] { # if_1 |
| $B2: { # true |
| %2:i32 = spirv.add<i32> 2i, 2i |
| exit_if # if_1 |
| } |
| $B3: { # false |
| exit_if # if_1 |
| } |
| } |
| if true [t: $B4, f: $B5] { # if_2 |
| $B4: { # true |
| %3:i32 = spirv.add<i32> 1i, 1i |
| ret |
| } |
| $B5: { # false |
| unreachable |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_Hoisting) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %main_start = OpLabel |
| %1 = OpCopyObject %i32 %one |
| OpSelectionMerge %bb_merge None |
| OpBranchConditional %true %bb_true %bb_false |
| %bb_true = OpLabel |
| %2 = OpCopyObject %i32 %two |
| OpBranch %bb_merge |
| %bb_false = OpLabel |
| OpReturn |
| %bb_merge = OpLabel |
| %3 = OpIAdd %i32 %2 %2 |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = let 1i |
| %3:i32 = if true [t: $B2, f: $B3] { # if_1 |
| $B2: { # true |
| %4:i32 = let 2i |
| exit_if %4 # if_1 |
| } |
| $B3: { # false |
| ret |
| } |
| } |
| %5:i32 = spirv.add<i32> %3, %3 |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_HoistingIntoNested) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %two = OpConstant %i32 2 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %502 = OpLabel |
| OpSelectionMerge %50 None |
| OpBranchConditional %true %20 %30 |
| %20 = OpLabel |
| %200 = OpCopyObject %i32 %two |
| OpBranch %50 |
| %30 = OpLabel |
| OpReturn |
| %50 = OpLabel |
| OpSelectionMerge %60 None |
| OpBranchConditional %true %70 %80 |
| %70 = OpLabel |
| %201 = OpIAdd %i32 %200 %200 |
| OpBranch %60 |
| %80 = OpLabel |
| OpReturn |
| %60 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = if true [t: $B2, f: $B3] { # if_1 |
| $B2: { # true |
| %3:i32 = let 2i |
| exit_if %3 # if_1 |
| } |
| $B3: { # false |
| ret |
| } |
| } |
| if true [t: $B4, f: $B5] { # if_2 |
| $B4: { # true |
| %4:i32 = spirv.add<i32> %2, %2 |
| exit_if # if_2 |
| } |
| $B5: { # false |
| ret |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_HoistingIntoParentNested) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %two = OpConstant %i32 2 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %1 = OpLabel |
| OpSelectionMerge %90 None |
| OpBranchConditional %true %10 %15 |
| %15 = OpLabel |
| OpReturn |
| %10 = OpLabel |
| OpSelectionMerge %50 None |
| OpBranchConditional %true %20 %30 |
| %20 = OpLabel |
| %200 = OpCopyObject %i32 %two |
| OpBranch %50 |
| %30 = OpLabel |
| OpReturn |
| %50 = OpLabel |
| OpSelectionMerge %60 None |
| OpBranchConditional %true %70 %80 |
| %70 = OpLabel |
| %201 = OpIAdd %i32 %200 %200 |
| OpBranch %60 |
| %80 = OpLabel |
| OpReturn |
| %60 = OpLabel |
| OpBranch %90 |
| %90 = OpLabel |
| %202 = OpIAdd %i32 %200 %200 |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = if true [t: $B2, f: $B3] { # if_1 |
| $B2: { # true |
| %3:i32 = if true [t: $B4, f: $B5] { # if_2 |
| $B4: { # true |
| %4:i32 = let 2i |
| exit_if %4 # if_2 |
| } |
| $B5: { # false |
| ret |
| } |
| } |
| if true [t: $B6, f: $B7] { # if_3 |
| $B6: { # true |
| %5:i32 = spirv.add<i32> %3, %3 |
| exit_if # if_3 |
| } |
| $B7: { # false |
| ret |
| } |
| } |
| exit_if %3 # if_1 |
| } |
| $B3: { # false |
| ret |
| } |
| } |
| %6:i32 = spirv.add<i32> %2, %2 |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_DuplicateTrue_UsedValue) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %main_start = OpLabel |
| %1 = OpCopyObject %i32 %one |
| OpBranchConditional %true %bb_true %bb_true |
| %bb_true = OpLabel |
| %2 = OpCopyObject %i32 %two |
| %22 = OpCopyObject %i32 %2 |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = let 1i |
| %3:bool = or true, true |
| if %3 [t: $B2, f: $B3] { # if_1 |
| $B2: { # true |
| %4:i32 = let 2i |
| %5:i32 = let %4 |
| ret |
| } |
| $B3: { # false |
| unreachable |
| } |
| } |
| unreachable |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_DuplicateTrue_Premerge) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %12 = OpLabel |
| OpSelectionMerge %13 None |
| OpBranchConditional %true %14 %15 |
| %14 = OpLabel |
| OpBranch %16 |
| %15 = OpLabel |
| OpBranch %16 |
| %16 = OpLabel |
| %1 = OpCopyObject %i32 %one |
| %2 = OpCopyObject %i32 %1 |
| OpBranch %13 |
| %13 = OpLabel |
| %3 = OpCopyObject %i32 %1 |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| if true [t: $B2, f: $B3] { # if_1 |
| $B2: { # true |
| exit_if # if_1 |
| } |
| $B3: { # false |
| exit_if # if_1 |
| } |
| } |
| %2:i32 = if true [t: $B4, f: $B5] { # if_2 |
| $B4: { # true |
| %3:i32 = let 1i |
| %4:i32 = let %3 |
| exit_if %3 # if_2 |
| } |
| $B5: { # false |
| unreachable |
| } |
| } |
| %5:i32 = let %2 |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Switch) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpSelectionMerge %99 None |
| OpSwitch %two %98 20 %20 |
| %20 = OpLabel |
| %21 = OpIAdd %i32 %one %two |
| OpBranch %99 |
| %98 = OpLabel |
| %22 = OpIAdd %i32 %two %two |
| OpBranch %99 |
| %99 = OpLabel |
| %23 = OpIAdd %i32 %three %two |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| switch 2i [c: (default, $B2), c: (20i, $B3)] { # switch_1 |
| $B2: { # case |
| %2:i32 = spirv.add<i32> 2i, 2i |
| exit_switch # switch_1 |
| } |
| $B3: { # case |
| %3:i32 = spirv.add<i32> 1i, 2i |
| exit_switch # switch_1 |
| } |
| } |
| %4:i32 = spirv.add<i32> 3i, 2i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Switch_DefaultIsMerge) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %two = OpConstant %i32 2 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpSelectionMerge %99 None |
| OpSwitch %two %99 20 %20 |
| %20 = OpLabel |
| %21 = OpIAdd %i32 %two %two |
| OpBranch %99 |
| %99 = OpLabel |
| %98 = OpIAdd %i32 %two %two |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| switch 2i [c: (default, $B2), c: (20i, $B3)] { # switch_1 |
| $B2: { # case |
| exit_switch # switch_1 |
| } |
| $B3: { # case |
| %2:i32 = spirv.add<i32> 2i, 2i |
| exit_switch # switch_1 |
| } |
| } |
| %3:i32 = spirv.add<i32> 2i, 2i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Switch_DefaultIsCase) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %two = OpConstant %i32 2 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpSelectionMerge %99 None |
| OpSwitch %two %20 20 %20 |
| %20 = OpLabel |
| %21 = OpIAdd %i32 %two %two |
| OpBranch %99 |
| %99 = OpLabel |
| %98 = OpIAdd %i32 %two %two |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| switch 2i [c: (20i default, $B2)] { # switch_1 |
| $B2: { # case |
| %2:i32 = spirv.add<i32> 2i, 2i |
| exit_switch # switch_1 |
| } |
| } |
| %3:i32 = spirv.add<i32> 2i, 2i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Switch_MuliselectorCase) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %two = OpConstant %i32 2 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpSelectionMerge %99 None |
| OpSwitch %two %50 20 %20 30 %20 |
| %20 = OpLabel |
| %21 = OpIAdd %i32 %two %two |
| OpBranch %99 |
| %50 = OpLabel |
| OpReturn |
| %99 = OpLabel |
| %98 = OpIAdd %i32 %two %two |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| switch 2i [c: (default, $B2), c: (20i 30i, $B3)] { # switch_1 |
| $B2: { # case |
| ret |
| } |
| $B3: { # case |
| %2:i32 = spirv.add<i32> 2i, 2i |
| exit_switch # switch_1 |
| } |
| } |
| %3:i32 = spirv.add<i32> 2i, 2i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Switch_OnlyDefault) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpSelectionMerge %99 None |
| OpSwitch %two %20 |
| %20 = OpLabel |
| %51 = OpIAdd %i32 %one %two |
| OpBranch %99 |
| %99 = OpLabel |
| %52 = OpIAdd %i32 %two %two |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| switch 2i [c: (default, $B2)] { # switch_1 |
| $B2: { # case |
| %2:i32 = spirv.add<i32> 1i, 2i |
| exit_switch # switch_1 |
| } |
| } |
| %3:i32 = spirv.add<i32> 2i, 2i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Switch_IfBreak) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %u32 = OpTypeInt 32 0 |
| %bool = OpTypeBool |
| %one = OpConstant %u32 1 |
| %two = OpConstant %u32 2 |
| %three = OpConstant %u32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpSelectionMerge %99 None |
| OpSwitch %two %50 20 %20 50 %50 |
| %99 = OpLabel |
| OpReturn |
| %20 = OpLabel |
| OpSelectionMerge %49 None |
| OpBranchConditional %true %99 %40 ; break-if |
| %40 = OpLabel |
| OpBranch %49 |
| %49 = OpLabel |
| %101 = OpIAdd %u32 %two %three |
| OpBranch %99 |
| %50 = OpLabel |
| OpSelectionMerge %79 None |
| OpBranchConditional %true %60 %99 ; break-unless |
| %60 = OpLabel |
| OpBranch %79 |
| %79 = OpLabel ; dominated by 60, so must follow 60 |
| %100 = OpIAdd %u32 %one %two |
| OpBranch %99 |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| switch 2u [c: (50u default, $B2), c: (20u, $B3)] { # switch_1 |
| $B2: { # case |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| exit_if # if_1 |
| } |
| $B5: { # false |
| exit_switch # switch_1 |
| } |
| } |
| %2:u32 = spirv.add<u32> 1u, 2u |
| exit_switch # switch_1 |
| } |
| $B3: { # case |
| if true [t: $B6, f: $B7] { # if_2 |
| $B6: { # true |
| exit_switch # switch_1 |
| } |
| $B7: { # false |
| exit_if # if_2 |
| } |
| } |
| %3:u32 = spirv.add<u32> 2u, 3u |
| exit_switch # switch_1 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Switch_Nest_If_In_Case) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %u32 = OpTypeInt 32 0 |
| %bool = OpTypeBool |
| %one = OpConstant %u32 1 |
| %two = OpConstant %u32 2 |
| %three = OpConstant %u32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpSelectionMerge %99 None |
| OpSwitch %two %50 20 %20 50 %50 |
| %99 = OpLabel |
| OpReturn |
| %20 = OpLabel |
| OpSelectionMerge %49 None |
| OpBranchConditional %true %30 %40 |
| %49 = OpLabel |
| OpBranch %99 |
| %30 = OpLabel |
| OpBranch %49 |
| %40 = OpLabel |
| OpBranch %49 |
| %50 = OpLabel |
| OpSelectionMerge %79 None |
| OpBranchConditional %false %60 %70 |
| %79 = OpLabel |
| OpBranch %99 |
| %60 = OpLabel |
| OpBranch %79 |
| %70 = OpLabel |
| OpBranch %79 |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| switch 2u [c: (50u default, $B2), c: (20u, $B3)] { # switch_1 |
| $B2: { # case |
| if false [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| exit_if # if_1 |
| } |
| $B5: { # false |
| exit_if # if_1 |
| } |
| } |
| exit_switch # switch_1 |
| } |
| $B3: { # case |
| if true [t: $B6, f: $B7] { # if_2 |
| $B6: { # true |
| exit_if # if_2 |
| } |
| $B7: { # false |
| exit_if # if_2 |
| } |
| } |
| exit_switch # switch_1 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Switch_HoistFromDefault) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %u32 = OpTypeInt 32 0 |
| %bool = OpTypeBool |
| %one = OpConstant %u32 1 |
| %two = OpConstant %u32 2 |
| %three = OpConstant %u32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpSelectionMerge %99 None |
| OpSwitch %two %20 |
| %20 = OpLabel |
| %100 = OpCopyObject %u32 %one |
| OpBranch %99 |
| %99 = OpLabel |
| %102 = OpIAdd %u32 %100 %100 |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:u32 = switch 2u [c: (default, $B2)] { # switch_1 |
| $B2: { # case |
| %3:u32 = let 1u |
| exit_switch %3 # switch_1 |
| } |
| } |
| %4:u32 = spirv.add<u32> %2, %2 |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Switch_HoistFromCase) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %u32 = OpTypeInt 32 0 |
| %bool = OpTypeBool |
| %one = OpConstant %u32 1 |
| %two = OpConstant %u32 2 |
| %three = OpConstant %u32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpSelectionMerge %99 None |
| OpSwitch %two %50 20 %20 |
| %20 = OpLabel |
| %100 = OpCopyObject %u32 %one |
| OpBranch %99 |
| %50 = OpLabel |
| OpReturn |
| %99 = OpLabel |
| %102 = OpIAdd %u32 %100 %100 |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:u32 = switch 2u [c: (default, $B2), c: (20u, $B3)] { # switch_1 |
| $B2: { # case |
| ret |
| } |
| $B3: { # case |
| %3:u32 = let 1u |
| exit_switch %3 # switch_1 |
| } |
| } |
| %4:u32 = spirv.add<u32> %2, %2 |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserDeathTest, Switch_Fallthrough) { |
| auto src = R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpSelectionMerge %99 None |
| OpSwitch %two %80 20 %20 30 %30 |
| %20 = OpLabel |
| %50 = OpIAdd %i32 %one %two |
| OpBranch %30 ; fall through |
| %30 = OpLabel |
| %51 = OpIAdd %i32 %two %two |
| OpBranch %99 |
| %80 = OpLabel |
| %52 = OpIAdd %i32 %three %two |
| OpBranch %99 |
| %99 = OpLabel |
| %53 = OpIAdd %i32 %one %three |
| OpReturn |
| OpFunctionEnd |
| )"; |
| EXPECT_DEATH_IF_SUPPORTED({ auto _ = Run(src); }, "internal compiler error"); |
| } |
| |
| TEST_F(SpirvParserTest, Switch_IfBreakInCase) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpSelectionMerge %99 None |
| OpSwitch %one %50 20 %20 50 %50 |
| %99 = OpLabel |
| OpReturn |
| %20 = OpLabel |
| OpSelectionMerge %49 None |
| OpBranchConditional %true %99 %40 ; break-if |
| %40 = OpLabel |
| OpBranch %49 |
| %49 = OpLabel |
| OpBranch %99 |
| %50 = OpLabel |
| OpSelectionMerge %79 None |
| OpBranchConditional %true %60 %99 ; break-unless |
| %60 = OpLabel |
| OpBranch %79 |
| %79 = OpLabel ; dominated by 60, so must follow 60 |
| OpBranch %99 |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| switch 1i [c: (50i default, $B2), c: (20i, $B3)] { # switch_1 |
| $B2: { # case |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| exit_if # if_1 |
| } |
| $B5: { # false |
| exit_switch # switch_1 |
| } |
| } |
| exit_switch # switch_1 |
| } |
| $B3: { # case |
| if true [t: $B6, f: $B7] { # if_2 |
| $B6: { # true |
| exit_switch # switch_1 |
| } |
| $B7: { # false |
| exit_if # if_2 |
| } |
| } |
| exit_switch # switch_1 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Switch_CaseCanBeMerge) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpSelectionMerge %99 None |
| OpSwitch %one %99 20 %99 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| switch 1i [c: (20i default, $B2)] { # switch_1 |
| $B2: { # case |
| exit_switch # switch_1 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| %1 = OpIAdd %i32 %one %two |
| OpLoopMerge %99 %20 None |
| OpBranchConditional %true %20 %99 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %2:i32 = spirv.add<i32> 1i, 2i |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| continue # -> $B3 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_Infinite_Branch) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %20 None |
| OpBranch %20 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_Infinite_BranchConditional) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %20 None |
| OpBranchConditional %true %20 %20 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %2:bool = or true, true |
| if %2 [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| continue # -> $B3 |
| } |
| $B5: { # false |
| unreachable |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_BranchConditional) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %40 None |
| OpBranchConditional %true %30 %99 |
| %30 = OpLabel |
| OpBranch %40 |
| %40 = OpLabel |
| OpBranch %20 ; back edge |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| continue # -> $B3 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_Branch) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %40 None |
| OpBranch %30 |
| %30 = OpLabel |
| OpBranchConditional %true %40 %99 |
| %40 = OpLabel |
| OpBranch %20 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| continue # -> $B3 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_BranchMergeOrContinue) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %30 = OpIAdd %i32 %one %two |
| OpBranch %20 |
| %20 = OpLabel |
| %31 = OpIAdd %i32 %two %three |
| OpLoopMerge %99 %20 None |
| OpBranchConditional %true %99 %20 |
| %99 = OpLabel |
| %32 = OpIAdd %i32 %two %three |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 2i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 3i |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| exit_loop # loop_1 |
| } |
| $B5: { # false |
| continue # -> $B3 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| %4:i32 = spirv.add<i32> 2i, 3i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_HeaderHasBreakIf) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %50 None |
| OpBranchConditional %true %30 %99 ; like While |
| %30 = OpLabel ; trivial body |
| OpBranch %50 |
| %50 = OpLabel |
| OpBranch %20 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| continue # -> $B3 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_HeaderHasBreakUnless) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %50 None |
| OpBranchConditional %true %99 %30 ; has break-unless |
| %30 = OpLabel ; trivial body |
| OpBranch %50 |
| %50 = OpLabel |
| OpBranch %20 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| exit_loop # loop_1 |
| } |
| $B5: { # false |
| continue # -> $B3 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_BodyHasBreak) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %50 None |
| OpBranchConditional %true %30 %99 |
| %30 = OpLabel |
| OpBranch %99 ; break |
| %50 = OpLabel |
| OpBranch %20 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| exit_loop # loop_1 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_BodyHasBreakIf) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %50 None |
| OpBranchConditional %true %30 %99 |
| %30 = OpLabel |
| OpBranchConditional %false %99 %40 ; break-if |
| %40 = OpLabel |
| OpBranch %50 |
| %50 = OpLabel |
| OpBranch %20 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| if false [t: $B6, f: $B7] { # if_2 |
| $B6: { # true |
| exit_loop # loop_1 |
| } |
| $B7: { # false |
| continue # -> $B3 |
| } |
| } |
| exit_if # if_1 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_BodyHasBreakUnless) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %50 None |
| OpBranchConditional %true %30 %99 |
| %30 = OpLabel |
| OpBranchConditional %false %40 %99 ; break-unless |
| %40 = OpLabel |
| OpBranch %50 |
| %50 = OpLabel |
| OpBranch %20 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| if false [t: $B6, f: $B7] { # if_2 |
| $B6: { # true |
| continue # -> $B3 |
| } |
| $B7: { # false |
| exit_loop # loop_1 |
| } |
| } |
| exit_if # if_1 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_BodyIf) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %50 None |
| OpBranchConditional %true %30 %99 |
| %30 = OpLabel |
| OpSelectionMerge %49 None |
| OpBranchConditional %false %40 %45 ; nested if |
| %40 = OpLabel |
| OpBranch %49 |
| %45 = OpLabel |
| OpBranch %49 |
| %49 = OpLabel |
| OpBranch %50 |
| %50 = OpLabel |
| OpBranch %20 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| if false [t: $B6, f: $B7] { # if_2 |
| $B6: { # true |
| exit_if # if_2 |
| } |
| $B7: { # false |
| exit_if # if_2 |
| } |
| } |
| continue # -> $B3 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_BodyIfBreak) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %50 None |
| OpBranchConditional %true %30 %99 |
| %30 = OpLabel |
| OpSelectionMerge %49 None |
| OpBranchConditional %false %40 %49 ; nested if |
| %40 = OpLabel |
| OpBranch %99 ; break from nested if |
| %49 = OpLabel |
| OpBranch %50 |
| %50 = OpLabel |
| OpBranch %20 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| if false [t: $B6, f: $B7] { # if_2 |
| $B6: { # true |
| exit_loop # loop_1 |
| } |
| $B7: { # false |
| exit_if # if_2 |
| } |
| } |
| continue # -> $B3 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_BodyHasContinueIf) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %50 None |
| OpBranchConditional %true %30 %99 |
| %30 = OpLabel |
| OpBranchConditional %false %50 %40 ; continue-if |
| %40 = OpLabel |
| OpBranch %50 |
| %50 = OpLabel |
| OpBranch %20 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| if false [t: $B6, f: $B7] { # if_2 |
| $B6: { # true |
| continue # -> $B3 |
| } |
| $B7: { # false |
| continue # -> $B3 |
| } |
| } |
| exit_if # if_1 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_BodyHasContinueUnless) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %50 None |
| OpBranchConditional %true %30 %99 |
| %30 = OpLabel |
| OpBranchConditional %false %40 %50 ; continue-unless |
| %40 = OpLabel |
| OpBranch %50 |
| %50 = OpLabel |
| OpBranch %20 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| if false [t: $B6, f: $B7] { # if_2 |
| $B6: { # true |
| continue # -> $B3 |
| } |
| $B7: { # false |
| continue # -> $B3 |
| } |
| } |
| exit_if # if_1 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_Body_If_Continue) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %50 None |
| OpBranchConditional %true %30 %99 |
| %30 = OpLabel |
| OpSelectionMerge %49 None |
| OpBranchConditional %false %40 %49 ; nested if |
| %40 = OpLabel |
| OpBranch %50 ; continue from nested if |
| %49 = OpLabel |
| OpBranch %50 |
| %50 = OpLabel |
| OpBranch %20 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| if false [t: $B6, f: $B7] { # if_2 |
| $B6: { # true |
| continue # -> $B3 |
| } |
| $B7: { # false |
| exit_if # if_2 |
| } |
| } |
| continue # -> $B3 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_ContinueUseBodyValue) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %50 None |
| OpBranchConditional %true %30 %99 |
| %30 = OpLabel |
| %40 = OpIAdd %i32 %one %two |
| OpBranch %50 |
| %50 = OpLabel |
| %41 = OpIAdd %i32 %40 %40 |
| OpBranch %20 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| %2:i32 = spirv.add<i32> 1i, 2i |
| continue %2 # -> $B3 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3 (%3:i32): { # continuing |
| %4:i32 = spirv.add<i32> %3, %3 |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_Body_Switch) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %50 None |
| OpBranchConditional %true %30 %99 |
| %30 = OpLabel |
| OpSelectionMerge %49 None |
| OpSwitch %one %49 40 %40 45 %45 ; fully nested switch |
| %40 = OpLabel |
| OpBranch %49 |
| %45 = OpLabel |
| OpBranch %49 |
| %49 = OpLabel |
| OpBranch %50 |
| %50 = OpLabel |
| OpBranch %20 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| switch 1i [c: (default, $B6), c: (40i, $B7), c: (45i, $B8)] { # switch_1 |
| $B6: { # case |
| exit_switch # switch_1 |
| } |
| $B7: { # case |
| exit_switch # switch_1 |
| } |
| $B8: { # case |
| exit_switch # switch_1 |
| } |
| } |
| continue # -> $B3 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_Body_Switch_CaseContinues) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %50 None |
| OpBranchConditional %true %30 %99 |
| %30 = OpLabel |
| OpSelectionMerge %49 None |
| OpSwitch %one %49 40 %40 45 %45 |
| %40 = OpLabel |
| OpBranch %50 ; continue bypasses switch merge |
| %45 = OpLabel |
| OpBranch %49 |
| %49 = OpLabel |
| OpBranch %50 |
| %50 = OpLabel |
| OpBranch %20 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| switch 1i [c: (default, $B6), c: (40i, $B7), c: (45i, $B8)] { # switch_1 |
| $B6: { # case |
| exit_switch # switch_1 |
| } |
| $B7: { # case |
| continue # -> $B3 |
| } |
| $B8: { # case |
| exit_switch # switch_1 |
| } |
| } |
| continue # -> $B3 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_Continue_Sequence) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %50 None |
| OpBranchConditional %true %30 %99 |
| %30 = OpLabel |
| OpBranch %50 |
| %50 = OpLabel |
| OpBranch %60 |
| %60 = OpLabel |
| OpBranch %20 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| continue # -> $B3 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_Continue_ContainsIf) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %50 None |
| OpBranchConditional %true %30 %99 |
| %30 = OpLabel |
| OpBranch %50 |
| %50 = OpLabel |
| OpSelectionMerge %89 None |
| OpBranchConditional %true %60 %70 |
| %89 = OpLabel |
| OpBranch %20 ; backedge |
| %60 = OpLabel |
| OpBranch %89 |
| %70 = OpLabel |
| OpBranch %89 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| continue # -> $B3 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| if true [t: $B6, f: $B7] { # if_2 |
| $B6: { # true |
| exit_if # if_2 |
| } |
| $B7: { # false |
| exit_if # if_2 |
| } |
| } |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_Continue_HasBreakIf) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %50 None |
| OpBranchConditional %true %30 %99 |
| %30 = OpLabel |
| OpBranch %50 |
| %50 = OpLabel |
| OpBranchConditional %false %99 %20 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| continue # -> $B3 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| break_if false # -> [t: exit_loop loop_1, f: $B2] |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_Continue_HasBreakUnless) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %50 None |
| OpBranchConditional %true %30 %99 |
| %30 = OpLabel |
| OpBranch %50 |
| %50 = OpLabel |
| OpBranchConditional %false %20 %99 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| continue # -> $B3 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| %2:bool = not false |
| break_if %2 # -> [t: exit_loop loop_1, f: $B2] |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_Loop) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %50 None |
| OpBranchConditional %true %30 %99 |
| %30 = OpLabel |
| OpLoopMerge %49 %40 None |
| OpBranchConditional %false %35 %49 |
| %35 = OpLabel |
| OpBranch %37 |
| %37 = OpLabel |
| OpBranch %40 |
| %40 = OpLabel ; inner loop's continue |
| OpBranch %30 ; backedge |
| %49 = OpLabel ; inner loop's merge |
| OpBranch %50 |
| %50 = OpLabel ; outer loop's continue |
| OpBranch %20 ; outer loop's backege |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| loop [b: $B6, c: $B7] { # loop_2 |
| $B6: { # body |
| if false [t: $B8, f: $B9] { # if_2 |
| $B8: { # true |
| continue # -> $B7 |
| } |
| $B9: { # false |
| exit_loop # loop_2 |
| } |
| } |
| unreachable |
| } |
| $B7: { # continuing |
| next_iteration # -> $B6 |
| } |
| } |
| continue # -> $B3 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_Loop_InnerBreak) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %50 None |
| OpBranchConditional %true %30 %99 |
| %30 = OpLabel |
| OpLoopMerge %49 %40 None |
| OpBranchConditional %false %35 %49 |
| %35 = OpLabel |
| OpBranchConditional %true %49 %37 ; break to inner merge |
| %37 = OpLabel |
| OpBranch %40 |
| %40 = OpLabel ; inner loop's continue |
| OpBranch %30 ; backedge |
| %49 = OpLabel ; inner loop's merge |
| OpBranch %50 |
| %50 = OpLabel ; outer loop's continue |
| OpBranch %20 ; outer loop's backege |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| loop [b: $B6, c: $B7] { # loop_2 |
| $B6: { # body |
| if false [t: $B8, f: $B9] { # if_2 |
| $B8: { # true |
| if true [t: $B10, f: $B11] { # if_3 |
| $B10: { # true |
| exit_loop # loop_2 |
| } |
| $B11: { # false |
| continue # -> $B7 |
| } |
| } |
| exit_if # if_2 |
| } |
| $B9: { # false |
| exit_loop # loop_2 |
| } |
| } |
| unreachable |
| } |
| $B7: { # continuing |
| next_iteration # -> $B6 |
| } |
| } |
| continue # -> $B3 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_Loop_InnerContinue) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %50 None |
| OpBranchConditional %true %30 %99 |
| %30 = OpLabel |
| OpLoopMerge %49 %40 None |
| OpBranchConditional %false %35 %49 |
| %35 = OpLabel |
| OpBranchConditional %true %37 %49 ; continue to inner continue target |
| %37 = OpLabel |
| OpBranch %40 |
| %40 = OpLabel ; inner loop's continue |
| OpBranch %30 ; backedge |
| %49 = OpLabel ; inner loop's merge |
| OpBranch %50 |
| %50 = OpLabel ; outer loop's continue |
| OpBranch %20 ; outer loop's backege |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| loop [b: $B6, c: $B7] { # loop_2 |
| $B6: { # body |
| if false [t: $B8, f: $B9] { # if_2 |
| $B8: { # true |
| if true [t: $B10, f: $B11] { # if_3 |
| $B10: { # true |
| continue # -> $B7 |
| } |
| $B11: { # false |
| exit_loop # loop_2 |
| } |
| } |
| exit_if # if_2 |
| } |
| $B9: { # false |
| exit_loop # loop_2 |
| } |
| } |
| unreachable |
| } |
| $B7: { # continuing |
| next_iteration # -> $B6 |
| } |
| } |
| continue # -> $B3 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_Loop_InnerContinueBreaks) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %50 None |
| OpBranchConditional %true %30 %99 |
| %30 = OpLabel |
| OpLoopMerge %49 %40 None |
| OpBranchConditional %false %35 %49 |
| %35 = OpLabel |
| OpBranch %37 |
| %37 = OpLabel |
| OpBranch %40 |
| %40 = OpLabel ; inner loop's continue |
| OpBranchConditional %true %30 %49 ; backedge and inner break |
| %49 = OpLabel ; inner loop's merge |
| OpBranch %50 |
| %50 = OpLabel ; outer loop's continue |
| OpBranch %20 ; outer loop's backege |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| loop [b: $B6, c: $B7] { # loop_2 |
| $B6: { # body |
| if false [t: $B8, f: $B9] { # if_2 |
| $B8: { # true |
| continue # -> $B7 |
| } |
| $B9: { # false |
| exit_loop # loop_2 |
| } |
| } |
| unreachable |
| } |
| $B7: { # continuing |
| %2:bool = not true |
| break_if %2 # -> [t: exit_loop loop_2, f: $B6] |
| } |
| } |
| continue # -> $B3 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_MergeBlockIsLoop) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpSelectionMerge %50 None |
| OpBranchConditional %true %20 %50 |
| %20 = OpLabel |
| OpBranch %50 |
| ; %50 is the merge block for the selection starting at 10, |
| ; and its own continue target. |
| %50 = OpLabel |
| OpLoopMerge %99 %50 None |
| OpBranchConditional %true %50 %99 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| if true [t: $B2, f: $B3] { # if_1 |
| $B2: { # true |
| exit_if # if_1 |
| } |
| $B3: { # false |
| exit_if # if_1 |
| } |
| } |
| loop [b: $B4, c: $B5] { # loop_1 |
| $B4: { # body |
| if true [t: $B6, f: $B7] { # if_2 |
| $B6: { # true |
| continue # -> $B5 |
| } |
| $B7: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B5: { # continuing |
| next_iteration # -> $B4 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, MergeIsAlsoMultiBlockLoopHeader) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpSelectionMerge %50 None |
| OpBranchConditional %true %20 %50 |
| %20 = OpLabel |
| OpBranch %50 |
| ; %50 is the merge block for the selection starting at 10, |
| ; and a loop block header but not its own continue target. |
| %50 = OpLabel |
| OpLoopMerge %99 %60 None |
| OpBranchConditional %true %60 %99 |
| %60 = OpLabel |
| OpBranch %50 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| if true [t: $B2, f: $B3] { # if_1 |
| $B2: { # true |
| exit_if # if_1 |
| } |
| $B3: { # false |
| exit_if # if_1 |
| } |
| } |
| loop [b: $B4, c: $B5] { # loop_1 |
| $B4: { # body |
| if true [t: $B6, f: $B7] { # if_2 |
| $B6: { # true |
| continue # -> $B5 |
| } |
| $B7: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B5: { # continuing |
| next_iteration # -> $B4 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| // Exercises the hard case where we a single OpBranchConditional has both |
| // IfBreak and Forward edges, within the true-branch clause. |
| TEST_F(SpirvParserTest, IfBreak_FromThen_ForwardWithinThen) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %two |
| OpSelectionMerge %99 None |
| OpBranchConditional %true %20 %50 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %three |
| OpBranchConditional %false %99 %30 ; kIfBreak with kForward |
| %30 = OpLabel ; still in then clause |
| %36 = OpIAdd %i32 %one %three |
| OpBranch %99 |
| %50 = OpLabel ; else clause |
| %37 = OpIAdd %i32 %three %one |
| OpBranch %99 |
| %99 = OpLabel |
| %38 = OpIAdd %i32 %three %two |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 2i |
| if true [t: $B2, f: $B3] { # if_1 |
| $B2: { # true |
| %3:i32 = spirv.add<i32> 2i, 3i |
| if false [t: $B4, f: $B5] { # if_2 |
| $B4: { # true |
| exit_if # if_2 |
| } |
| $B5: { # false |
| %4:i32 = spirv.add<i32> 1i, 3i |
| exit_if # if_2 |
| } |
| } |
| exit_if # if_1 |
| } |
| $B3: { # false |
| %5:i32 = spirv.add<i32> 3i, 1i |
| exit_if # if_1 |
| } |
| } |
| %6:i32 = spirv.add<i32> 3i, 2i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| // Exercises the hard case where we a single OpBranchConditional has both |
| // IfBreak and Forward edges, within the false-branch clause. |
| TEST_F(SpirvParserTest, IfBreak_FromElse_ForwardWithinElse) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %35 = OpIAdd %i32 %one %one |
| OpSelectionMerge %99 None |
| OpBranchConditional %true %20 %50 |
| %20 = OpLabel |
| %36 = OpIAdd %i32 %two %two |
| OpBranch %99 |
| %50 = OpLabel ; else clause |
| %37 = OpIAdd %i32 %three %three |
| OpBranchConditional %false %99 %80 ; kIfBreak with kForward |
| %80 = OpLabel ; still in then clause |
| %38 = OpIAdd %i32 %one %two |
| OpBranch %99 |
| %99 = OpLabel |
| %39 = OpIAdd %i32 %two %three |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| if true [t: $B2, f: $B3] { # if_1 |
| $B2: { # true |
| %3:i32 = spirv.add<i32> 2i, 2i |
| exit_if # if_1 |
| } |
| $B3: { # false |
| %4:i32 = spirv.add<i32> 3i, 3i |
| if false [t: $B4, f: $B5] { # if_2 |
| $B4: { # true |
| exit_if # if_2 |
| } |
| $B5: { # false |
| %5:i32 = spirv.add<i32> 1i, 2i |
| exit_if # if_2 |
| } |
| } |
| exit_if # if_1 |
| } |
| } |
| %6:i32 = spirv.add<i32> 2i, 3i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| // This is a combination of the previous two, but also adding a premerge. |
| // We have IfBreak and Forward edges from the same OpBranchConditional, and |
| // this occurs in the true-branch clause, the false-branch clause, and within |
| // the premerge clause. Flow guards have to be sprinkled in lots of places. |
| TEST_F(SpirvParserTest, IfBreak_FromThenAndElseWithForward_Premerge) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %35 = OpIAdd %i32 %one %one |
| OpSelectionMerge %99 None |
| OpBranchConditional %true %20 %50 |
| %20 = OpLabel ; then |
| %36 = OpIAdd %i32 %two %two |
| OpBranchConditional %false %21 %99 ; kForward and kIfBreak |
| %21 = OpLabel ; still in then clause |
| %37 = OpIAdd %i32 %three %three |
| OpBranch %80 ; to premerge |
| %50 = OpLabel ; else clause |
| %38 = OpIAdd %i32 %one %two |
| OpBranchConditional %false %99 %51 ; kIfBreak with kForward |
| %51 = OpLabel ; still in else clause |
| %39 = OpIAdd %i32 %two %three |
| OpBranch %80 ; to premerge |
| %80 = OpLabel ; premerge |
| %41 = OpIAdd %i32 %three %two |
| OpBranchConditional %true %81 %99 |
| %81 = OpLabel ; premerge |
| %42 = OpIAdd %i32 %three %one |
| OpBranch %99 |
| %99 = OpLabel |
| %43 = OpIAdd %i32 %two %one |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| %execute_premerge:ptr<function, bool, read_write> = var true |
| if true [t: $B2, f: $B3] { # if_1 |
| $B2: { # true |
| %4:i32 = spirv.add<i32> 2i, 2i |
| if false [t: $B4, f: $B5] { # if_2 |
| $B4: { # true |
| %5:i32 = spirv.add<i32> 3i, 3i |
| exit_if # if_2 |
| } |
| $B5: { # false |
| store %execute_premerge, false |
| exit_if # if_2 |
| } |
| } |
| exit_if # if_1 |
| } |
| $B3: { # false |
| %6:i32 = spirv.add<i32> 1i, 2i |
| if false [t: $B6, f: $B7] { # if_3 |
| $B6: { # true |
| store %execute_premerge, false |
| exit_if # if_3 |
| } |
| $B7: { # false |
| %7:i32 = spirv.add<i32> 2i, 3i |
| exit_if # if_3 |
| } |
| } |
| exit_if # if_1 |
| } |
| } |
| %8:bool = load %execute_premerge |
| if %8 [t: $B8, f: $B9] { # if_4 |
| $B8: { # true |
| %9:i32 = spirv.add<i32> 3i, 2i |
| if true [t: $B10, f: $B11] { # if_5 |
| $B10: { # true |
| %10:i32 = spirv.add<i32> 3i, 1i |
| exit_if # if_5 |
| } |
| $B11: { # false |
| exit_if # if_5 |
| } |
| } |
| exit_if # if_4 |
| } |
| $B9: { # false |
| unreachable |
| } |
| } |
| %11:i32 = spirv.add<i32> 2i, 1i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_SingleBlock_BothBackedge) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %35 = OpIAdd %i32 %one %one |
| OpBranch %20 |
| %20 = OpLabel |
| %36 = OpIAdd %i32 %two %two |
| OpLoopMerge %99 %20 None |
| OpBranchConditional %true %20 %20 |
| %99 = OpLabel |
| %37 = OpIAdd %i32 %three %three |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 2i |
| %4:bool = or true, true |
| if %4 [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| continue # -> $B3 |
| } |
| $B5: { # false |
| unreachable |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| %5:i32 = spirv.add<i32> 3i, 3i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_SingleBlock_UnconditionalBackege) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %35 = OpIAdd %i32 %one %one |
| OpBranch %20 |
| %20 = OpLabel |
| %36 = OpIAdd %i32 %two %two |
| OpLoopMerge %99 %20 None |
| OpBranch %20 |
| %99 = OpLabel |
| %37 = OpIAdd %i32 %three %three |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 2i |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| %4:i32 = spirv.add<i32> 3i, 3i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_Unconditional_Body_SingleBlockContinue) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %35 = OpIAdd %i32 %one %one |
| OpBranch %20 |
| %20 = OpLabel |
| %36 = OpIAdd %i32 %two %two |
| OpLoopMerge %99 %50 None |
| OpBranch %30 |
| %30 = OpLabel |
| %37 = OpIAdd %i32 %three %three |
| OpBranch %50 |
| %50 = OpLabel |
| %38 = OpIAdd %i32 %one %two |
| OpBranch %20 |
| %99 = OpLabel |
| %39 = OpIAdd %i32 %two %three |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 2i |
| %4:i32 = spirv.add<i32> 3i, 3i |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| %5:i32 = spirv.add<i32> 1i, 2i |
| next_iteration # -> $B2 |
| } |
| } |
| %6:i32 = spirv.add<i32> 2i, 3i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_Unconditional_Body_MultiBlockContinue) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpLoopMerge %99 %50 None |
| OpBranch %30 |
| %30 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| OpBranch %50 |
| %50 = OpLabel |
| %37 = OpIAdd %i32 %one %three |
| OpBranch %60 |
| %60 = OpLabel |
| %38 = OpIAdd %i32 %one %two |
| OpBranch %20 |
| %99 = OpLabel |
| %39 = OpIAdd %i32 %two %three |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 2i |
| %4:i32 = spirv.add<i32> 3i, 3i |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| %5:i32 = spirv.add<i32> 1i, 3i |
| %6:i32 = spirv.add<i32> 1i, 2i |
| next_iteration # -> $B2 |
| } |
| } |
| %7:i32 = spirv.add<i32> 2i, 3i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_Unconditional_Body_ContinueNestIf) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpLoopMerge %99 %50 None |
| OpBranch %30 |
| %30 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| OpBranch %50 |
| %50 = OpLabel ; continue target; also if-header |
| %37 = OpIAdd %i32 %one %two |
| OpSelectionMerge %80 None |
| OpBranchConditional %false %60 %80 |
| %60 = OpLabel |
| %38 = OpIAdd %i32 %two %three |
| OpBranch %80 |
| %80 = OpLabel |
| %39 = OpIAdd %i32 %one %three |
| OpBranch %20 |
| %99 = OpLabel |
| %41 = OpIAdd %i32 %three %two |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 2i |
| %4:i32 = spirv.add<i32> 3i, 3i |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| %5:i32 = spirv.add<i32> 1i, 2i |
| if false [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| %6:i32 = spirv.add<i32> 2i, 3i |
| exit_if # if_1 |
| } |
| $B5: { # false |
| exit_if # if_1 |
| } |
| } |
| %7:i32 = spirv.add<i32> 1i, 3i |
| next_iteration # -> $B2 |
| } |
| } |
| %8:i32 = spirv.add<i32> 3i, 2i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| // Test case where both branches exit. e.g both go to merge. |
| TEST_F(SpirvParserTest, Loop_Never) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpLoopMerge %99 %80 None |
| OpBranchConditional %true %99 %99 |
| %80 = OpLabel ; continue target |
| %35 = OpIAdd %i32 %two %two |
| OpBranch %20 |
| %99 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %2:i32 = spirv.add<i32> 1i, 1i |
| %3:bool = or true, true |
| if %3 [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| exit_loop # loop_1 |
| } |
| $B5: { # false |
| unreachable |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| %4:i32 = spirv.add<i32> 2i, 2i |
| next_iteration # -> $B2 |
| } |
| } |
| %5:i32 = spirv.add<i32> 3i, 3i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_TrueToBody_FalseBreaks) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpLoopMerge %99 %80 None |
| OpBranchConditional %true %30 %99 |
| %30 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpBranch %80 |
| %80 = OpLabel ; continue target |
| %36 = OpIAdd %i32 %three %three |
| OpBranch %20 |
| %99 = OpLabel |
| %37 = OpIAdd %i32 %one %two |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %2:i32 = spirv.add<i32> 1i, 1i |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| %3:i32 = spirv.add<i32> 2i, 2i |
| continue # -> $B3 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| %4:i32 = spirv.add<i32> 3i, 3i |
| next_iteration # -> $B2 |
| } |
| } |
| %5:i32 = spirv.add<i32> 1i, 2i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_FalseToBody_TrueBreaks) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpLoopMerge %99 %80 None |
| OpBranchConditional %true %30 %99 |
| %30 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpBranch %80 |
| %80 = OpLabel ; continue target |
| %36 = OpIAdd %i32 %three %three |
| OpBranch %20 |
| %99 = OpLabel |
| %37 = OpIAdd %i32 %one %three |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %2:i32 = spirv.add<i32> 1i, 1i |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| %3:i32 = spirv.add<i32> 2i, 2i |
| continue # -> $B3 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| %4:i32 = spirv.add<i32> 3i, 3i |
| next_iteration # -> $B2 |
| } |
| } |
| %5:i32 = spirv.add<i32> 1i, 3i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| // By construction, it has to come from nested code. |
| TEST_F(SpirvParserTest, Loop_NestedIfContinue) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %80 None |
| OpBranch %30 |
| %30 = OpLabel |
| OpSelectionMerge %50 None |
| OpBranchConditional %true %40 %50 |
| %40 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %80 ; continue edge |
| %50 = OpLabel ; inner selection merge |
| %35 = OpIAdd %i32 %two %two |
| OpBranch %80 |
| %80 = OpLabel ; continue target |
| %36 = OpIAdd %i32 %three %three |
| OpBranch %20 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| %2:i32 = spirv.add<i32> 1i, 1i |
| continue # -> $B3 |
| } |
| $B5: { # false |
| exit_if # if_1 |
| } |
| } |
| %3:i32 = spirv.add<i32> 2i, 2i |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| %4:i32 = spirv.add<i32> 3i, 3i |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_BodyAlwaysBreaks) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %80 None |
| OpBranch %30 |
| %30 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %99 ; break is here |
| %80 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpBranch %20 ; backedge |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %2:i32 = spirv.add<i32> 1i, 1i |
| exit_loop # loop_1 |
| } |
| $B3: { # continuing |
| %3:i32 = spirv.add<i32> 2i, 2i |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| // The else-branch has a continue but it's skipped because it's from a |
| // block that immediately precedes the continue construct. |
| TEST_F(SpirvParserTest, Loop_BodyConditionallyBreaks_FromTrue) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %80 None |
| OpBranch %30 |
| %30 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranchConditional %true %99 %80 |
| %80 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpBranch %20 ; backedge |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %2:i32 = spirv.add<i32> 1i, 1i |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| exit_loop # loop_1 |
| } |
| $B5: { # false |
| continue # -> $B3 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| %3:i32 = spirv.add<i32> 2i, 2i |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| // The else-branch has a continue but it's skipped because it's from a |
| // block that immediately precedes the continue construct. |
| TEST_F(SpirvParserTest, Loop_BodyConditionallyBreaks_FromFalse) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %80 None |
| OpBranch %30 |
| %30 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranchConditional %true %80 %99 |
| %80 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpBranch %20 ; backedge |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %2:i32 = spirv.add<i32> 1i, 1i |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| continue # -> $B3 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| %3:i32 = spirv.add<i32> 2i, 2i |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_BodyConditionallyBreaks_FromTrue_Early) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %80 None |
| OpBranch %30 |
| %30 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranchConditional %true %99 %70 |
| %70 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpBranch %80 |
| %80 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| OpBranch %20 ; backedge |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %2:i32 = spirv.add<i32> 1i, 1i |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| exit_loop # loop_1 |
| } |
| $B5: { # false |
| %3:i32 = spirv.add<i32> 2i, 2i |
| continue # -> $B3 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| %4:i32 = spirv.add<i32> 3i, 3i |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Loop_BodyConditionallyBreaks_FromFalse_Early) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %80 None |
| OpBranch %30 |
| %30 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranchConditional %true %70 %99 |
| %70 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpBranch %80 |
| %80 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| OpBranch %20 ; backedge |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %2:i32 = spirv.add<i32> 1i, 1i |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| %3:i32 = spirv.add<i32> 2i, 2i |
| continue # -> $B3 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| %4:i32 = spirv.add<i32> 3i, 3i |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Switch_Case_SintValue) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpSelectionMerge %99 None |
| ; SPIR-V assembler doesn't support negative literals in switch |
| OpSwitch %one %99 20 %20 2000000000 %30 !4000000000 %40 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %one %one |
| OpBranch %99 |
| %30 = OpLabel |
| %36 = OpIAdd %i32 %two %two |
| OpBranch %99 |
| %40 = OpLabel |
| %37 = OpIAdd %i32 %three %three |
| OpBranch %99 |
| %99 = OpLabel |
| %38 = OpIAdd %i32 %one %three |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| switch 1i [c: (default, $B2), c: (20i, $B3), c: (2000000000i, $B4), c: (-294967296i, $B5)] { # switch_1 |
| $B2: { # case |
| exit_switch # switch_1 |
| } |
| $B3: { # case |
| %3:i32 = spirv.add<i32> 1i, 1i |
| exit_switch # switch_1 |
| } |
| $B4: { # case |
| %4:i32 = spirv.add<i32> 2i, 2i |
| exit_switch # switch_1 |
| } |
| $B5: { # case |
| %5:i32 = spirv.add<i32> 3i, 3i |
| exit_switch # switch_1 |
| } |
| } |
| %6:i32 = spirv.add<i32> 1i, 3i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Switch_Case_UintValue) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %u32 = OpTypeInt 32 0 |
| %bool = OpTypeBool |
| %one = OpConstant %u32 1 |
| %two = OpConstant %u32 2 |
| %three = OpConstant %u32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %u32 %one %one |
| OpSelectionMerge %99 None |
| OpSwitch %one %99 20 %20 2000000000 %30 50 %40 |
| %20 = OpLabel |
| %35 = OpIAdd %u32 %two %two |
| OpBranch %99 |
| %30 = OpLabel |
| %36 = OpIAdd %u32 %three %three |
| OpBranch %99 |
| %40 = OpLabel |
| %37 = OpIAdd %u32 %one %two |
| OpBranch %99 |
| %99 = OpLabel |
| %38 = OpIAdd %u32 %two %three |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:u32 = spirv.add<u32> 1u, 1u |
| switch 1u [c: (default, $B2), c: (20u, $B3), c: (2000000000u, $B4), c: (50u, $B5)] { # switch_1 |
| $B2: { # case |
| exit_switch # switch_1 |
| } |
| $B3: { # case |
| %3:u32 = spirv.add<u32> 2u, 2u |
| exit_switch # switch_1 |
| } |
| $B4: { # case |
| %4:u32 = spirv.add<u32> 3u, 3u |
| exit_switch # switch_1 |
| } |
| $B5: { # case |
| %5:u32 = spirv.add<u32> 1u, 2u |
| exit_switch # switch_1 |
| } |
| } |
| %6:u32 = spirv.add<u32> 2u, 3u |
| ret |
| } |
| } |
| )"); |
| } |
| |
| // When the break is not last in its case, we must emit a 'break' |
| TEST_F(SpirvParserTest, Branch_SwitchBreak_NotLastInCase) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpSelectionMerge %99 None |
| OpSwitch %one %99 20 %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpSelectionMerge %50 None |
| OpBranchConditional %true %40 %50 |
| %40 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| OpBranch %99 ; branch to merge. Not last in case |
| %50 = OpLabel ; inner merge |
| %37 = OpIAdd %i32 %one %two |
| OpBranch %99 |
| %99 = OpLabel |
| %38 = OpIAdd %i32 %two %three |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| switch 1i [c: (default, $B2), c: (20i, $B3)] { # switch_1 |
| $B2: { # case |
| exit_switch # switch_1 |
| } |
| $B3: { # case |
| %3:i32 = spirv.add<i32> 2i, 2i |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| %4:i32 = spirv.add<i32> 3i, 3i |
| exit_switch # switch_1 |
| } |
| $B5: { # false |
| exit_if # if_1 |
| } |
| } |
| %5:i32 = spirv.add<i32> 1i, 2i |
| exit_switch # switch_1 |
| } |
| } |
| %6:i32 = spirv.add<i32> 2i, 3i |
| ret |
| } |
| } |
| )"); |
| |
| // auto* expect = R"(var_1 = 1u; |
| // switch(42u) { |
| // case 20u: { |
| // var_1 = 20u; |
| // if (false) { |
| // var_1 = 40u; |
| // break; |
| // } |
| // var_1 = 50u; |
| // } |
| // default: { |
| // } |
| // } |
| // var_1 = 7u; |
| // return; |
| } |
| |
| TEST_F(SpirvParserTest, Branch_LoopBreak_MultiBlockLoop_FromBody) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %80 None |
| OpBranch %30 |
| %30 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %99 ; break is here |
| %80 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpBranch %20 ; backedge |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %2:i32 = spirv.add<i32> 1i, 1i |
| exit_loop # loop_1 |
| } |
| $B3: { # continuing |
| %3:i32 = spirv.add<i32> 2i, 2i |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, |
| Branch_LoopBreak_MultiBlockLoop_FromContinueConstructEnd_Conditional_BreakIf) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %80 None |
| OpBranch %80 |
| %80 = OpLabel ; continue target |
| %34 = OpIAdd %i32 %one %one |
| OpBranchConditional %true %99 %20 ; exit, and backedge |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| %2:i32 = spirv.add<i32> 1i, 1i |
| break_if true # -> [t: exit_loop loop_1, f: $B2] |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Branch_LoopBreak_MultiBlockLoop_FromContinueConstructEnd_Conditional) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %80 None |
| OpBranch %80 |
| %80 = OpLabel ; continue target |
| %34 = OpIAdd %i32 %one %one |
| OpBranchConditional %true %20 %99 ; backedge, and exit |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| %2:i32 = spirv.add<i32> 1i, 1i |
| %3:bool = not true |
| break_if %3 # -> [t: exit_loop loop_1, f: $B2] |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Branch_LoopBreak_FromContinueConstructTail) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %50 None |
| OpBranchConditional %true %30 %99 |
| %30 = OpLabel |
| OpBranch %50 |
| %50 = OpLabel |
| OpBranch %60 |
| %60 = OpLabel |
| OpBranchConditional %true %20 %99 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| continue # -> $B3 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| %2:bool = not true |
| break_if %2 # -> [t: exit_loop loop_1, f: $B2] |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Branch_LoopContinue_LastInLoopConstruct) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %80 None |
| OpBranch %30 |
| %30 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %80 ; continue edge from last block before continue target |
| %80 = OpLabel ; continue target |
| %35 = OpIAdd %i32 %two %two |
| OpBranch %20 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %2:i32 = spirv.add<i32> 1i, 1i |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| %3:i32 = spirv.add<i32> 2i, 2i |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| // By construction, it has to come from nested code. |
| TEST_F(SpirvParserTest, Branch_LoopContinue_BeforeLast) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpBranch %20 |
| %20 = OpLabel |
| OpLoopMerge %99 %80 None |
| OpBranch %30 |
| %30 = OpLabel |
| OpSelectionMerge %50 None |
| OpBranchConditional %true %40 %50 |
| %40 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %80 ; continue edge |
| %50 = OpLabel ; inner selection merge |
| %35 = OpIAdd %i32 %two %two |
| OpBranch %80 |
| %80 = OpLabel ; continue target |
| %36 = OpIAdd %i32 %three %three |
| OpBranch %20 |
| %99 = OpLabel |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| %2:i32 = spirv.add<i32> 1i, 1i |
| continue # -> $B3 |
| } |
| $B5: { # false |
| exit_if # if_1 |
| } |
| } |
| %3:i32 = spirv.add<i32> 2i, 2i |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| %4:i32 = spirv.add<i32> 3i, 3i |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, Branch_LoopContinue_FromSwitch) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpLoopMerge %99 %80 None |
| OpBranch %30 |
| %30 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| OpSelectionMerge %79 None |
| OpSwitch %one %79 40 %40 |
| %40 = OpLabel |
| %37 = OpIAdd %i32 %one %two |
| OpBranch %80 ; continue edge |
| %79 = OpLabel ; switch merge |
| %38 = OpIAdd %i32 %three %two |
| OpBranch %80 |
| %80 = OpLabel ; continue target |
| %39 = OpIAdd %i32 %two %three |
| OpBranch %20 |
| %99 = OpLabel |
| %41 = OpIAdd %i32 %three %one |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 2i |
| %4:i32 = spirv.add<i32> 3i, 3i |
| switch 1i [c: (default, $B4), c: (40i, $B5)] { # switch_1 |
| $B4: { # case |
| exit_switch # switch_1 |
| } |
| $B5: { # case |
| %5:i32 = spirv.add<i32> 1i, 2i |
| continue # -> $B3 |
| } |
| } |
| %6:i32 = spirv.add<i32> 3i, 2i |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| %7:i32 = spirv.add<i32> 2i, 3i |
| next_iteration # -> $B2 |
| } |
| } |
| %8:i32 = spirv.add<i32> 3i, 1i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| // When unconditional, the if-break must be last in the then clause. |
| TEST_F(SpirvParserTest, Branch_IfBreak_FromThen) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpSelectionMerge %99 None |
| OpBranchConditional %true %30 %99 |
| %30 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %99 |
| %99 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| if true [t: $B2, f: $B3] { # if_1 |
| $B2: { # true |
| %2:i32 = spirv.add<i32> 1i, 1i |
| exit_if # if_1 |
| } |
| $B3: { # false |
| exit_if # if_1 |
| } |
| } |
| %3:i32 = spirv.add<i32> 2i, 2i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| // When unconditional, the if-break must be last in the else clause. |
| TEST_F(SpirvParserTest, Branch_IfBreak_FromElse) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpSelectionMerge %99 None |
| OpBranchConditional %true %99 %30 |
| %30 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %99 |
| %99 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| if true [t: $B2, f: $B3] { # if_1 |
| $B2: { # true |
| exit_if # if_1 |
| } |
| $B3: { # false |
| %2:i32 = spirv.add<i32> 1i, 1i |
| exit_if # if_1 |
| } |
| } |
| %3:i32 = spirv.add<i32> 2i, 2i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_Back_SingleBlock_LoopBreak_OnTrue) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpLoopMerge %99 %20 None |
| OpBranchConditional %true %99 %20 |
| %99 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 2i |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| exit_loop # loop_1 |
| } |
| $B5: { # false |
| continue # -> $B3 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| %4:i32 = spirv.add<i32> 3i, 3i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_Back_SingleBlock_LoopBreak_OnFalse) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpLoopMerge %99 %20 None |
| OpBranchConditional %true %20 %99 |
| %99 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 2i |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| continue # -> $B3 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| %4:i32 = spirv.add<i32> 3i, 3i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_Back_MultiBlock_LoopBreak_OnTrue) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpLoopMerge %99 %80 None |
| OpBranch %80 |
| %80 = OpLabel |
| OpBranchConditional %true %99 %20 |
| %99 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 2i |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| break_if true # -> [t: exit_loop loop_1, f: $B2] |
| } |
| } |
| %4:i32 = spirv.add<i32> 3i, 3i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_Back_MultiBlock_LoopBreak_OnFalse) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpLoopMerge %99 %80 None |
| OpBranch %80 |
| %80 = OpLabel |
| OpBranchConditional %true %20 %99 |
| %99 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 2i |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| %4:bool = not true |
| break_if %4 # -> [t: exit_loop loop_1, f: $B2] |
| } |
| } |
| %5:i32 = spirv.add<i32> 3i, 3i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| // When the break is not last in its case, we must emit a 'break' |
| TEST_F(SpirvParserTest, BranchConditional_SwitchBreak_SwitchBreak_NotLastInCase) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpSelectionMerge %99 None |
| OpSwitch %one %99 20 %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpSelectionMerge %50 None |
| OpBranchConditional %true %40 %50 |
| %40 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| OpBranchConditional %false %99 %99 ; branch to merge. Not last in case |
| %50 = OpLabel ; inner merge |
| %37 = OpIAdd %i32 %one %two |
| OpBranch %99 |
| %99 = OpLabel |
| %38 = OpIAdd %i32 %one %three |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| switch 1i [c: (default, $B2), c: (20i, $B3)] { # switch_1 |
| $B2: { # case |
| exit_switch # switch_1 |
| } |
| $B3: { # case |
| %3:i32 = spirv.add<i32> 2i, 2i |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| %4:i32 = spirv.add<i32> 3i, 3i |
| %5:bool = or false, true |
| if %5 [t: $B6, f: $B7] { # if_2 |
| $B6: { # true |
| exit_switch # switch_1 |
| } |
| $B7: { # false |
| unreachable |
| } |
| } |
| exit_if # if_1 |
| } |
| $B5: { # false |
| exit_if # if_1 |
| } |
| } |
| %6:i32 = spirv.add<i32> 1i, 2i |
| exit_switch # switch_1 |
| } |
| } |
| %7:i32 = spirv.add<i32> 1i, 3i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_SwitchBreak_Continue_OnTrue) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpLoopMerge %99 %80 None |
| OpBranch %30 |
| %30 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| OpSelectionMerge %79 None |
| OpSwitch %one %79 40 %40 |
| %40 = OpLabel |
| %37 = OpIAdd %i32 %one %two |
| OpBranchConditional %true %80 %79 ; break; continue on true |
| %79 = OpLabel |
| %38 = OpIAdd %i32 %one %three |
| OpBranch %80 |
| %80 = OpLabel ; continue target |
| %39 = OpIAdd %i32 %two %three |
| OpBranch %20 |
| %99 = OpLabel ; loop merge |
| %41 = OpIAdd %i32 %three %one |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 2i |
| %4:i32 = spirv.add<i32> 3i, 3i |
| switch 1i [c: (default, $B4), c: (40i, $B5)] { # switch_1 |
| $B4: { # case |
| exit_switch # switch_1 |
| } |
| $B5: { # case |
| %5:i32 = spirv.add<i32> 1i, 2i |
| if true [t: $B6, f: $B7] { # if_1 |
| $B6: { # true |
| continue # -> $B3 |
| } |
| $B7: { # false |
| exit_switch # switch_1 |
| } |
| } |
| exit_switch # switch_1 |
| } |
| } |
| %6:i32 = spirv.add<i32> 1i, 3i |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| %7:i32 = spirv.add<i32> 2i, 3i |
| next_iteration # -> $B2 |
| } |
| } |
| %8:i32 = spirv.add<i32> 3i, 1i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_SwitchBreak_Continue_OnFalse) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpLoopMerge %99 %80 None |
| OpBranch %30 |
| %30 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| OpSelectionMerge %79 None |
| OpSwitch %one %79 40 %40 |
| %40 = OpLabel |
| %37 = OpIAdd %i32 %one %two |
| OpBranchConditional %true %79 %80 ; break; continue on false |
| %79 = OpLabel |
| %38 = OpIAdd %i32 %one %three |
| OpBranch %80 |
| %80 = OpLabel ; continue target |
| %39 = OpIAdd %i32 %two %three |
| OpBranch %20 |
| %99 = OpLabel ; loop merge |
| %41 = OpIAdd %i32 %three %one |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 2i |
| %4:i32 = spirv.add<i32> 3i, 3i |
| switch 1i [c: (default, $B4), c: (40i, $B5)] { # switch_1 |
| $B4: { # case |
| exit_switch # switch_1 |
| } |
| $B5: { # case |
| %5:i32 = spirv.add<i32> 1i, 2i |
| if true [t: $B6, f: $B7] { # if_1 |
| $B6: { # true |
| exit_switch # switch_1 |
| } |
| $B7: { # false |
| continue # -> $B3 |
| } |
| } |
| exit_switch # switch_1 |
| } |
| } |
| %6:i32 = spirv.add<i32> 1i, 3i |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| %7:i32 = spirv.add<i32> 2i, 3i |
| next_iteration # -> $B2 |
| } |
| } |
| %8:i32 = spirv.add<i32> 3i, 1i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_SwitchBreak_Forward_OnTrue) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpSelectionMerge %99 None |
| OpSwitch %one %99 20 %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpBranchConditional %true %30 %99 ; break; forward on true |
| %30 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| OpBranch %99 |
| %99 = OpLabel ; switch merge |
| %37 = OpIAdd %i32 %one %two |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| switch 1i [c: (default, $B2), c: (20i, $B3)] { # switch_1 |
| $B2: { # case |
| exit_switch # switch_1 |
| } |
| $B3: { # case |
| %3:i32 = spirv.add<i32> 2i, 2i |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| %4:i32 = spirv.add<i32> 3i, 3i |
| exit_switch # switch_1 |
| } |
| $B5: { # false |
| exit_switch # switch_1 |
| } |
| } |
| exit_switch # switch_1 |
| } |
| } |
| %5:i32 = spirv.add<i32> 1i, 2i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_SwitchBreak_Forward_OnFalse) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpSelectionMerge %99 None |
| OpSwitch %one %99 20 %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpBranchConditional %true %99 %30 ; break; forward on false |
| %30 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| OpBranch %99 |
| %99 = OpLabel ; switch merge |
| %37 = OpIAdd %i32 %one %two |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| switch 1i [c: (default, $B2), c: (20i, $B3)] { # switch_1 |
| $B2: { # case |
| exit_switch # switch_1 |
| } |
| $B3: { # case |
| %3:i32 = spirv.add<i32> 2i, 2i |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| exit_switch # switch_1 |
| } |
| $B5: { # false |
| %4:i32 = spirv.add<i32> 3i, 3i |
| exit_switch # switch_1 |
| } |
| } |
| exit_switch # switch_1 |
| } |
| } |
| %5:i32 = spirv.add<i32> 1i, 2i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_LoopBreak_SingleBlock_LoopBreak) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpLoopMerge %99 %80 None |
| OpBranchConditional %true %99 %99 |
| %80 = OpLabel ; continue target |
| %36 = OpIAdd %i32 %three %three |
| OpBranch %20 |
| %99 = OpLabel |
| %37 = OpIAdd %i32 %one %three |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 2i |
| %4:bool = or true, true |
| if %4 [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| exit_loop # loop_1 |
| } |
| $B5: { # false |
| unreachable |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| %5:i32 = spirv.add<i32> 3i, 3i |
| next_iteration # -> $B2 |
| } |
| } |
| %6:i32 = spirv.add<i32> 1i, 3i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_LoopBreak_MultiBlock_LoopBreak) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpLoopMerge %99 %80 None |
| OpBranch %30 |
| %30 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| OpBranchConditional %true %99 %99 |
| %80 = OpLabel ; continue target |
| %37 = OpIAdd %i32 %one %two |
| OpBranch %20 |
| %99 = OpLabel |
| %38 = OpIAdd %i32 %one %three |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 2i |
| %4:i32 = spirv.add<i32> 3i, 3i |
| %5:bool = or true, true |
| if %5 [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| exit_loop # loop_1 |
| } |
| $B5: { # false |
| unreachable |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| %6:i32 = spirv.add<i32> 1i, 2i |
| next_iteration # -> $B2 |
| } |
| } |
| %7:i32 = spirv.add<i32> 1i, 3i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_LoopBreak_Continue_OnTrue) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpLoopMerge %99 %80 None |
| OpBranch %25 |
| ; Need this extra selection to make another block between |
| ; %30 and the continue target, so we actually induce a Continue |
| ; statement to exist. |
| %25 = OpLabel |
| OpSelectionMerge %40 None |
| OpBranchConditional %false %30 %40 |
| %30 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| ; break; continue on true |
| OpBranchConditional %false %80 %99 |
| %40 = OpLabel |
| %37 = OpIAdd %i32 %one %two |
| OpBranch %80 |
| %80 = OpLabel ; continue target |
| %38 = OpIAdd %i32 %one %three |
| OpBranch %20 |
| %99 = OpLabel |
| %39 = OpIAdd %i32 %two %three |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 2i |
| if false [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| %4:i32 = spirv.add<i32> 3i, 3i |
| if false [t: $B6, f: $B7] { # if_2 |
| $B6: { # true |
| continue # -> $B3 |
| } |
| $B7: { # false |
| exit_loop # loop_1 |
| } |
| } |
| exit_if # if_1 |
| } |
| $B5: { # false |
| exit_if # if_1 |
| } |
| } |
| %5:i32 = spirv.add<i32> 1i, 2i |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| %6:i32 = spirv.add<i32> 1i, 3i |
| next_iteration # -> $B2 |
| } |
| } |
| %7:i32 = spirv.add<i32> 2i, 3i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_LoopBreak_Continue_OnFalse) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpLoopMerge %99 %80 None |
| OpBranch %25 |
| ; Need this extra selection to make another block between |
| ; %30 and the continue target, so we actually induce a Continue |
| ; statement to exist. |
| %25 = OpLabel |
| OpSelectionMerge %40 None |
| OpBranchConditional %false %30 %40 |
| %30 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| ; break; continue on false |
| OpBranchConditional %true %99 %80 |
| %40 = OpLabel |
| %37 = OpIAdd %i32 %one %two |
| OpBranch %80 |
| %80 = OpLabel ; continue target |
| %38 = OpIAdd %i32 %one %three |
| OpBranch %20 |
| %99 = OpLabel |
| %39 = OpIAdd %i32 %two %three |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 2i |
| if false [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| %4:i32 = spirv.add<i32> 3i, 3i |
| if true [t: $B6, f: $B7] { # if_2 |
| $B6: { # true |
| exit_loop # loop_1 |
| } |
| $B7: { # false |
| continue # -> $B3 |
| } |
| } |
| exit_if # if_1 |
| } |
| $B5: { # false |
| exit_if # if_1 |
| } |
| } |
| %5:i32 = spirv.add<i32> 1i, 2i |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| %6:i32 = spirv.add<i32> 1i, 3i |
| next_iteration # -> $B2 |
| } |
| } |
| %7:i32 = spirv.add<i32> 2i, 3i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_LoopBreak_Forward_OnTrue) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpLoopMerge %99 %80 None |
| OpBranch %30 |
| %30 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| ; break; forward on true |
| OpBranchConditional %true %40 %99 |
| %40 = OpLabel |
| %37 = OpIAdd %i32 %one %two |
| OpBranch %80 |
| %80 = OpLabel ; continue target |
| %38 = OpIAdd %i32 %one %three |
| OpBranch %20 |
| %99 = OpLabel |
| %39 = OpIAdd %i32 %two %three |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 2i |
| %4:i32 = spirv.add<i32> 3i, 3i |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| %5:i32 = spirv.add<i32> 1i, 2i |
| continue # -> $B3 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| %6:i32 = spirv.add<i32> 1i, 3i |
| next_iteration # -> $B2 |
| } |
| } |
| %7:i32 = spirv.add<i32> 2i, 3i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_LoopBreak_Forward_OnFalse) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpLoopMerge %99 %80 None |
| OpBranch %30 |
| %30 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| ; break; forward on false |
| OpBranchConditional %true %99 %40 |
| %40 = OpLabel |
| %37 = OpIAdd %i32 %one %two |
| OpBranch %80 |
| %80 = OpLabel ; continue target |
| %38 = OpIAdd %i32 %one %three |
| OpBranch %20 |
| %99 = OpLabel |
| %39 = OpIAdd %i32 %two %three |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 2i |
| %4:i32 = spirv.add<i32> 3i, 3i |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| exit_loop # loop_1 |
| } |
| $B5: { # false |
| %5:i32 = spirv.add<i32> 1i, 2i |
| continue # -> $B3 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| %6:i32 = spirv.add<i32> 1i, 3i |
| next_iteration # -> $B2 |
| } |
| } |
| %7:i32 = spirv.add<i32> 2i, 3i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_Continue_Continue_FromHeader) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpLoopMerge %99 %80 None |
| OpBranchConditional %true %80 %80 ; to continue |
| %80 = OpLabel ; continue target |
| %36 = OpIAdd %i32 %three %three |
| OpBranch %20 |
| %99 = OpLabel |
| %37 = OpIAdd %i32 %one %two |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 2i |
| %4:bool = or true, true |
| if %4 [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| continue # -> $B3 |
| } |
| $B5: { # false |
| unreachable |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| %5:i32 = spirv.add<i32> 3i, 3i |
| next_iteration # -> $B2 |
| } |
| } |
| %6:i32 = spirv.add<i32> 1i, 2i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_Continue_Continue_AfterHeader_Unconditional) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpLoopMerge %99 %80 None |
| OpBranch %30 |
| %30 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| OpBranchConditional %true %80 %80 ; to continue |
| %80 = OpLabel ; continue target |
| %37 = OpIAdd %i32 %one %two |
| OpBranch %20 |
| %99 = OpLabel |
| %38 = OpIAdd %i32 %one %three |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 2i |
| %4:i32 = spirv.add<i32> 3i, 3i |
| %5:bool = or true, true |
| if %5 [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| continue # -> $B3 |
| } |
| $B5: { # false |
| unreachable |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| %6:i32 = spirv.add<i32> 1i, 2i |
| next_iteration # -> $B2 |
| } |
| } |
| %7:i32 = spirv.add<i32> 1i, 3i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| // Create an intervening block so we actually require a "continue" statement |
| TEST_F(SpirvParserTest, BranchConditional_Continue_Continue_AfterHeader_Conditional) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpLoopMerge %99 %80 None |
| OpBranch %30 |
| %30 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| OpSelectionMerge %50 None |
| OpBranchConditional %false %40 %50 |
| %40 = OpLabel |
| %37 = OpIAdd %i32 %one %two |
| OpBranchConditional %true %80 %80 ; to continue |
| %50 = OpLabel ; merge for selection |
| %38 = OpIAdd %i32 %one %three |
| OpBranch %80 |
| %80 = OpLabel ; continue target |
| %39 = OpIAdd %i32 %two %three |
| OpBranch %20 |
| %99 = OpLabel |
| %41 = OpIAdd %i32 %three %one |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 2i |
| %4:i32 = spirv.add<i32> 3i, 3i |
| if false [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| %5:i32 = spirv.add<i32> 1i, 2i |
| %6:bool = or true, true |
| if %6 [t: $B6, f: $B7] { # if_2 |
| $B6: { # true |
| continue # -> $B3 |
| } |
| $B7: { # false |
| unreachable |
| } |
| } |
| exit_if # if_1 |
| } |
| $B5: { # false |
| exit_if # if_1 |
| } |
| } |
| %7:i32 = spirv.add<i32> 1i, 3i |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| %8:i32 = spirv.add<i32> 2i, 3i |
| next_iteration # -> $B2 |
| } |
| } |
| %9:i32 = spirv.add<i32> 3i, 1i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| // Like the previous tests, but with an empty continuing clause. |
| TEST_F(SpirvParserTest, |
| BranchConditional_Continue_Continue_AfterHeader_Conditional_EmptyContinuing) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpLoopMerge %99 %80 None |
| OpBranch %30 |
| %30 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| OpSelectionMerge %50 None |
| OpBranchConditional %false %40 %50 |
| %40 = OpLabel |
| %37 = OpIAdd %i32 %one %two |
| OpBranchConditional %false %80 %80 ; to continue |
| %50 = OpLabel ; merge for selection |
| %38 = OpIAdd %i32 %one %three |
| OpBranch %80 |
| %80 = OpLabel ; continue target |
| ; no statements here. |
| OpBranch %20 |
| %99 = OpLabel |
| %39 = OpIAdd %i32 %two %three |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 2i |
| %4:i32 = spirv.add<i32> 3i, 3i |
| if false [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| %5:i32 = spirv.add<i32> 1i, 2i |
| %6:bool = or false, true |
| if %6 [t: $B6, f: $B7] { # if_2 |
| $B6: { # true |
| continue # -> $B3 |
| } |
| $B7: { # false |
| unreachable |
| } |
| } |
| exit_if # if_1 |
| } |
| $B5: { # false |
| exit_if # if_1 |
| } |
| } |
| %7:i32 = spirv.add<i32> 1i, 3i |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| %8:i32 = spirv.add<i32> 2i, 3i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_LoopContinue_FromSwitch) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpLoopMerge %99 %80 None |
| OpBranch %30 |
| %30 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| OpSelectionMerge %79 None |
| OpSwitch %one %79 40 %40 |
| %40 = OpLabel |
| %37 = OpIAdd %i32 %one %two |
| OpBranchConditional %false %80 %80; dup continue edge |
| %79 = OpLabel ; switch merge |
| %38 = OpIAdd %i32 %one %three |
| OpBranch %80 |
| %80 = OpLabel ; continue target |
| %39 = OpIAdd %i32 %two %three |
| OpBranch %20 |
| %99 = OpLabel |
| %41 = OpIAdd %i32 %three %one |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 2i |
| %4:i32 = spirv.add<i32> 3i, 3i |
| switch 1i [c: (default, $B4), c: (40i, $B5)] { # switch_1 |
| $B4: { # case |
| exit_switch # switch_1 |
| } |
| $B5: { # case |
| %5:i32 = spirv.add<i32> 1i, 2i |
| %6:bool = or false, true |
| if %6 [t: $B6, f: $B7] { # if_1 |
| $B6: { # true |
| continue # -> $B3 |
| } |
| $B7: { # false |
| unreachable |
| } |
| } |
| exit_switch # switch_1 |
| } |
| } |
| %7:i32 = spirv.add<i32> 1i, 3i |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| %8:i32 = spirv.add<i32> 2i, 3i |
| next_iteration # -> $B2 |
| } |
| } |
| %9:i32 = spirv.add<i32> 3i, 1i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_Continue_IfBreak_OnTrue) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpLoopMerge %99 %80 None |
| OpBranch %30 |
| %30 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| OpSelectionMerge %50 None |
| OpBranchConditional %false %40 %50 |
| %40 = OpLabel |
| %37 = OpIAdd %i32 %one %two |
| ; true to if's merge; false to continue |
| OpBranchConditional %true %50 %80 |
| %50 = OpLabel ; merge for selection |
| %38 = OpIAdd %i32 %one %three |
| OpBranch %80 |
| %80 = OpLabel ; continue target |
| %39 = OpIAdd %i32 %two %one |
| OpBranch %20 |
| %99 = OpLabel |
| %41 = OpIAdd %i32 %three %one |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 2i |
| %4:i32 = spirv.add<i32> 3i, 3i |
| if false [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| %5:i32 = spirv.add<i32> 1i, 2i |
| if true [t: $B6, f: $B7] { # if_2 |
| $B6: { # true |
| exit_if # if_2 |
| } |
| $B7: { # false |
| continue # -> $B3 |
| } |
| } |
| exit_if # if_1 |
| } |
| $B5: { # false |
| exit_if # if_1 |
| } |
| } |
| %6:i32 = spirv.add<i32> 1i, 3i |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| %7:i32 = spirv.add<i32> 2i, 1i |
| next_iteration # -> $B2 |
| } |
| } |
| %8:i32 = spirv.add<i32> 3i, 1i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_Continue_IfBreak_OnFalse) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %false = OpConstantFalse %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpLoopMerge %99 %80 None |
| OpBranch %30 |
| %30 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| OpSelectionMerge %50 None |
| OpBranchConditional %false %40 %50 |
| %40 = OpLabel |
| %37 = OpIAdd %i32 %one %two |
| ; false to if's merge; true to continue |
| OpBranchConditional %true %80 %50 |
| %50 = OpLabel ; merge for selection |
| %38 = OpIAdd %i32 %one %two |
| OpBranch %80 |
| %80 = OpLabel ; continue target |
| %39 = OpIAdd %i32 %two %three |
| OpBranch %20 |
| %99 = OpLabel |
| %41 = OpIAdd %i32 %three %one |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 2i |
| %4:i32 = spirv.add<i32> 3i, 3i |
| if false [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| %5:i32 = spirv.add<i32> 1i, 2i |
| if true [t: $B6, f: $B7] { # if_2 |
| $B6: { # true |
| continue # -> $B3 |
| } |
| $B7: { # false |
| exit_if # if_2 |
| } |
| } |
| exit_if # if_1 |
| } |
| $B5: { # false |
| exit_if # if_1 |
| } |
| } |
| %6:i32 = spirv.add<i32> 1i, 2i |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| %7:i32 = spirv.add<i32> 2i, 3i |
| next_iteration # -> $B2 |
| } |
| } |
| %8:i32 = spirv.add<i32> 3i, 1i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_Continue_Forward_OnTrue) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpLoopMerge %99 %80 None |
| OpBranch %30 |
| %30 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| ; continue; forward on true |
| OpBranchConditional %true %40 %80 |
| %40 = OpLabel |
| %37 = OpIAdd %i32 %one %two |
| OpBranch %80 |
| %80 = OpLabel ; continue target |
| %38 = OpIAdd %i32 %one %three |
| OpBranch %20 |
| %99 = OpLabel |
| %39 = OpIAdd %i32 %two %one |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 2i |
| %4:i32 = spirv.add<i32> 3i, 3i |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| %5:i32 = spirv.add<i32> 1i, 2i |
| continue # -> $B3 |
| } |
| $B5: { # false |
| continue # -> $B3 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| %6:i32 = spirv.add<i32> 1i, 3i |
| next_iteration # -> $B2 |
| } |
| } |
| %7:i32 = spirv.add<i32> 2i, 1i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_Continue_Forward_OnFalse) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranch %20 |
| %20 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpLoopMerge %99 %80 None |
| OpBranch %30 |
| %30 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| ; continue; forward on true |
| OpBranchConditional %true %80 %40 |
| %40 = OpLabel |
| %37 = OpIAdd %i32 %one %two |
| OpBranch %80 |
| %80 = OpLabel ; continue target |
| %38 = OpIAdd %i32 %one %three |
| OpBranch %20 |
| %99 = OpLabel |
| %39 = OpIAdd %i32 %two %one |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %3:i32 = spirv.add<i32> 2i, 2i |
| %4:i32 = spirv.add<i32> 3i, 3i |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| continue # -> $B3 |
| } |
| $B5: { # false |
| %5:i32 = spirv.add<i32> 1i, 2i |
| continue # -> $B3 |
| } |
| } |
| unreachable |
| } |
| $B3: { # continuing |
| %6:i32 = spirv.add<i32> 1i, 3i |
| next_iteration # -> $B2 |
| } |
| } |
| %7:i32 = spirv.add<i32> 2i, 1i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_IfBreak_IfBreak_Same) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpSelectionMerge %99 None |
| OpBranchConditional %true %99 %99 |
| %20 = OpLabel ; dead |
| %35 = OpIAdd %i32 %two %two |
| OpBranch %99 |
| %99 = OpLabel |
| %36 = OpIAdd %i32 %three %three |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| %3:bool = or true, true |
| if %3 [t: $B2, f: $B3] { # if_1 |
| $B2: { # true |
| exit_if # if_1 |
| } |
| $B3: { # false |
| unreachable |
| } |
| } |
| %4:i32 = spirv.add<i32> 3i, 3i |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(SpirvParserTest, BranchConditional_Forward_Forward_Same) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| %34 = OpIAdd %i32 %one %one |
| OpBranchConditional %true %99 %99; forward |
| %99 = OpLabel |
| %35 = OpIAdd %i32 %two %two |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| %2:i32 = spirv.add<i32> 1i, 1i |
| %3:bool = or true, true |
| if %3 [t: $B2, f: $B3] { # if_1 |
| $B2: { # true |
| %4:i32 = spirv.add<i32> 2i, 2i |
| ret |
| } |
| $B3: { # false |
| unreachable |
| } |
| } |
| unreachable |
| } |
| } |
| )"); |
| } |
| |
| // crbug.com/tint/243 |
| TEST_F(SpirvParserTest, IfSelection_TrueBranch_LoopBreak) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %5 = OpLabel |
| OpBranch %10 |
| %10 = OpLabel |
| OpLoopMerge %99 %90 None |
| OpBranch %20 |
| %20 = OpLabel |
| OpSelectionMerge %40 None |
| OpBranchConditional %true %99 %30 ; true branch breaking is ok |
| %30 = OpLabel |
| OpBranch %40 |
| %40 = OpLabel ; selection merge |
| OpBranch %90 |
| %90 = OpLabel ; continue target |
| OpBranch %10 ; backedge |
| %99 = OpLabel ; loop merge |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| exit_loop # loop_1 |
| } |
| $B5: { # false |
| exit_if # if_1 |
| } |
| } |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| // crbug.com/tint/243 |
| TEST_F(SpirvParserTest, TrueBranch_LoopContinue) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %5 = OpLabel |
| OpBranch %10 |
| %10 = OpLabel |
| OpLoopMerge %99 %90 None |
| OpBranch %20 |
| %20 = OpLabel |
| OpSelectionMerge %40 None |
| OpBranchConditional %true %90 %30 ; true branch continue is ok |
| %30 = OpLabel |
| OpBranch %40 |
| %40 = OpLabel ; selection merge |
| OpBranch %90 |
| %90 = OpLabel ; continue target |
| OpBranch %10 ; backedge |
| %99 = OpLabel ; loop merge |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| continue # -> $B3 |
| } |
| $B5: { # false |
| exit_if # if_1 |
| } |
| } |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| // crbug.com/tint/243 |
| TEST_F(SpirvParserTest, TrueBranch_SwitchBreak) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpSelectionMerge %99 None |
| OpSwitch %one %99 20 %20 |
| %20 = OpLabel |
| OpSelectionMerge %40 None |
| OpBranchConditional %true %99 %30 ; true branch switch break is ok |
| %30 = OpLabel |
| OpBranch %40 |
| %40 = OpLabel ; if-selection merge |
| OpBranch %99 |
| %99 = OpLabel ; switch merge |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| switch 1i [c: (default, $B2), c: (20i, $B3)] { # switch_1 |
| $B2: { # case |
| exit_switch # switch_1 |
| } |
| $B3: { # case |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| exit_switch # switch_1 |
| } |
| $B5: { # false |
| exit_if # if_1 |
| } |
| } |
| exit_switch # switch_1 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| // crbug.com/tint/243 |
| TEST_F(SpirvParserTest, FalseBranch_LoopBreak) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %5 = OpLabel |
| OpBranch %10 |
| %10 = OpLabel |
| OpLoopMerge %99 %90 None |
| OpBranch %20 |
| %20 = OpLabel |
| OpSelectionMerge %40 None |
| OpBranchConditional %true %30 %99 ; false branch breaking is ok |
| %30 = OpLabel |
| OpBranch %40 |
| %40 = OpLabel ; selection merge |
| OpBranch %90 |
| %90 = OpLabel ; continue target |
| OpBranch %10 ; backedge |
| %99 = OpLabel ; loop merge |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| exit_if # if_1 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| // crbug.com/tint/243 |
| TEST_F(SpirvParserTest, FalseBranch_LoopContinue) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %5 = OpLabel |
| OpBranch %10 |
| %10 = OpLabel |
| OpLoopMerge %99 %90 None |
| OpBranch %20 |
| %20 = OpLabel |
| OpSelectionMerge %40 None |
| OpBranchConditional %true %30 %90 ; false branch continue is ok |
| %30 = OpLabel |
| OpBranch %40 |
| %40 = OpLabel ; selection merge |
| OpBranch %90 |
| %90 = OpLabel ; continue target |
| OpBranch %10 ; backedge |
| %99 = OpLabel ; loop merge |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| exit_if # if_1 |
| } |
| $B5: { # false |
| continue # -> $B3 |
| } |
| } |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| // crbug.com/tint/243 |
| TEST_F(SpirvParserTest, FalseBranch_SwitchBreak) { |
| EXPECT_IR(R"( |
| OpCapability Shader |
| OpMemoryModel Logical GLSL450 |
| OpEntryPoint GLCompute %main "main" |
| OpExecutionMode %main LocalSize 1 1 1 |
| %void = OpTypeVoid |
| %i32 = OpTypeInt 32 1 |
| %bool = OpTypeBool |
| %one = OpConstant %i32 1 |
| %two = OpConstant %i32 2 |
| %three = OpConstant %i32 3 |
| %true = OpConstantTrue %bool |
| %ep_type = OpTypeFunction %void |
| %main = OpFunction %void None %ep_type |
| %10 = OpLabel |
| OpSelectionMerge %99 None |
| OpSwitch %one %99 20 %20 |
| %20 = OpLabel |
| OpSelectionMerge %40 None |
| OpBranchConditional %true %30 %99 ; false branch switch break is ok |
| %30 = OpLabel |
| OpBranch %40 |
| %40 = OpLabel ; if-selection merge |
| OpBranch %99 |
| %99 = OpLabel ; switch merge |
| OpReturn |
| OpFunctionEnd |
| )", |
| R"( |
| %main = @compute @workgroup_size(1u, 1u, 1u) func():void { |
| $B1: { |
| switch 1i [c: (default, $B2), c: (20i, $B3)] { # switch_1 |
| $B2: { # case |
| exit_switch # switch_1 |
| } |
| $B3: { # case |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| exit_if # if_1 |
| } |
| $B5: { # false |
| exit_switch # switch_1 |
| } |
| } |
| exit_switch # switch_1 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| } // namespace |
| } // namespace tint::spirv::reader |
