// 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, Phi_FromBlock) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %bool = OpTypeBool
        %int = OpTypeInt 32 1
       %true = OpConstantTrue %bool
      %int_2 = OpConstant %int 2
    %ep_type = OpTypeFunction %void
       %main = OpFunction %void None %ep_type
 %main_start = OpLabel
               OpBranch %1
          %1 = OpLabel
          %2 = OpPhi %int %int_2 %main_start
          %3 = OpIAdd %int %2 %2
               OpReturn
               OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    %2:i32 = spirv.add<i32> 2i, 2i
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, Phi_If_Undef) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %bool = OpTypeBool
        %int = OpTypeInt 32 1
       %true = OpConstantTrue %bool
    %ep_type = OpTypeFunction %void
       %main = OpFunction %void None %ep_type
 %main_start = OpLabel
               OpBranch %1
          %1 = OpLabel
               OpSelectionMerge %5 None
               OpBranchConditional %true %3 %4
          %3 = OpLabel
          %6 = OpUndef %int
               OpBranch %5
          %4 = OpLabel
          %7 = OpUndef %int
               OpBranch %5
          %5 = OpLabel
          %8 = OpPhi %int %6 %3 %7 %4
          %9 = OpIAdd %int %8 %8
               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
        exit_if 0i  # if_1
      }
      $B3: {  # false
        exit_if 0i  # if_1
      }
    }
    %3:i32 = spirv.add<i32> %2, %2
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, Phi_If_ThenAndElse) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %bool = OpTypeBool
        %int = OpTypeInt 32 1
       %true = OpConstantTrue %bool
      %int_1 = OpConstant %int 1
      %int_2 = OpConstant %int 2
      %int_3 = OpConstant %int 3
      %int_4 = OpConstant %int 4
    %ep_type = OpTypeFunction %void
       %main = OpFunction %void None %ep_type
 %main_start = OpLabel
               OpBranch %1
          %1 = OpLabel
               OpSelectionMerge %5 None
               OpBranchConditional %true %3 %4
          %3 = OpLabel
          %6 = OpIAdd %int %int_1 %int_2
               OpBranch %5
          %4 = OpLabel
          %7 = OpIAdd %int %int_3 %int_4
               OpBranch %5
          %5 = OpLabel
          %8 = OpPhi %int %6 %3 %7 %4
          %9 = OpIAdd %int %8 %8
               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 = spirv.add<i32> 1i, 2i
        exit_if %3  # if_1
      }
      $B3: {  # false
        %4:i32 = spirv.add<i32> 3i, 4i
        exit_if %4  # if_1
      }
    }
    %5:i32 = spirv.add<i32> %2, %2
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, Phi_If_ThenNoElse) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %bool = OpTypeBool
        %int = OpTypeInt 32 1
       %true = OpConstantTrue %bool
      %int_1 = OpConstant %int 1
      %int_2 = OpConstant %int 2
      %int_3 = OpConstant %int 3
      %int_4 = OpConstant %int 4
    %ep_type = OpTypeFunction %void
       %main = OpFunction %void None %ep_type
 %main_start = OpLabel
               OpBranch %1
          %1 = OpLabel
               OpSelectionMerge %5 None
               OpBranchConditional %true %3 %5
          %3 = OpLabel
          %6 = OpIAdd %int %int_1 %int_2
               OpBranch %5
          %5 = OpLabel
          %8 = OpPhi %int %int_2 %1 %6 %3
          %9 = OpIAdd %int %8 %8
               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 = spirv.add<i32> 1i, 2i
        exit_if %3  # if_1
      }
      $B3: {  # false
        exit_if 2i  # if_1
      }
    }
    %4:i32 = spirv.add<i32> %2, %2
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, Phi_If_NoThenElse) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %bool = OpTypeBool
        %int = OpTypeInt 32 1
       %true = OpConstantTrue %bool
      %int_1 = OpConstant %int 1
      %int_2 = OpConstant %int 2
      %int_3 = OpConstant %int 3
      %int_4 = OpConstant %int 4
    %ep_type = OpTypeFunction %void
       %main = OpFunction %void None %ep_type
 %main_start = OpLabel
               OpBranch %1
          %1 = OpLabel
               OpSelectionMerge %5 None
               OpBranchConditional %true %5 %3
          %3 = OpLabel
          %6 = OpIAdd %int %int_1 %int_2
               OpBranch %5
          %5 = OpLabel
          %8 = OpPhi %int %6 %3 %int_2 %1
          %9 = OpIAdd %int %8 %8
               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
        exit_if 2i  # if_1
      }
      $B3: {  # false
        %3:i32 = spirv.add<i32> 1i, 2i
        exit_if %3  # if_1
      }
    }
    %4:i32 = spirv.add<i32> %2, %2
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, Phi_Switch) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
          %3 = OpTypeFunction %void
        %int = OpTypeInt 32 1
      %int_0 = OpConstant %int 0
      %int_1 = OpConstant %int 1
      %int_2 = OpConstant %int 2
       %main = OpFunction %void None %3
          %4 = OpLabel
               OpSelectionMerge %13 None
               OpSwitch %int_0 %10 0 %11 1 %12
         %10 = OpLabel
               OpBranch %13
         %11 = OpLabel
               OpBranch %13
         %12 = OpLabel
               OpBranch %13
         %13 = OpLabel
         %14 = OpPhi %int %int_0 %10 %int_1 %11 %int_2 %12
         %15 = OpIAdd %int %14 %14
               OpReturn
               OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    %2:i32 = switch 0i [c: (default, $B2), c: (0i, $B3), c: (1i, $B4)] {  # switch_1
      $B2: {  # case
        exit_switch 0i  # switch_1
      }
      $B3: {  # case
        exit_switch 1i  # switch_1
      }
      $B4: {  # case
        exit_switch 2i  # switch_1
      }
    }
    %3:i32 = spirv.add<i32> %2, %2
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, Phi_Loop_ContinueIsHeader) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %bool = OpTypeBool
       %true = OpConstantTrue %bool
      %false = OpConstantFalse %bool
    %ep_type = OpTypeFunction %void
       %main = OpFunction %void None %ep_type
         %10 = OpLabel
               OpBranch %20
         %20 = OpLabel
        %101 = OpPhi %bool %true %10 %false %20
               OpLoopMerge %99 %20 None
               OpBranchConditional %101 %99 %20
         %99 = OpLabel
               OpReturn
               OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    loop [i: $B2, b: $B3, c: $B4] {  # loop_1
      $B2: {  # initializer
        next_iteration true  # -> $B3
      }
      $B3 (%2:bool): {  # body
        if %2 [t: $B5, f: $B6] {  # if_1
          $B5: {  # true
            exit_loop  # loop_1
          }
          $B6: {  # false
            continue  # -> $B4
          }
        }
        unreachable
      }
      $B4: {  # continuing
        next_iteration false  # -> $B3
      }
    }
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, Phi_Loop_WithContinue) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %bool = OpTypeBool
       %true = OpConstantTrue %bool
      %false = OpConstantFalse %bool
    %ep_type = OpTypeFunction %void
       %main = OpFunction %void None %ep_type
         %10 = OpLabel
               OpBranch %20
         %20 = OpLabel
        %101 = OpPhi %bool %true %10 %false %30
               OpLoopMerge %99 %30 None
               OpBranchConditional %101 %99 %25
         %25 = OpLabel
               OpBranch %30
         %30 = OpLabel
        %102 = OpCopyObject %bool %101
               OpBranch %20
         %99 = OpLabel
               OpReturn
               OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    loop [i: $B2, b: $B3, c: $B4] {  # loop_1
      $B2: {  # initializer
        next_iteration true  # -> $B3
      }
      $B3 (%2:bool): {  # body
        if %2 [t: $B5, f: $B6] {  # if_1
          $B5: {  # true
            exit_loop  # loop_1
          }
          $B6: {  # false
            continue  # -> $B4
          }
        }
        unreachable
      }
      $B4: {  # continuing
        %3:bool = let %2
        next_iteration false  # -> $B3
      }
    }
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, Phi_Loop_WithContinue_PhiInContinue) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %bool = OpTypeBool
       %true = OpConstantTrue %bool
      %false = OpConstantFalse %bool
    %ep_type = OpTypeFunction %void
       %main = OpFunction %void None %ep_type
         %10 = OpLabel
               OpBranch %20
         %20 = OpLabel
               OpLoopMerge %99 %30 None
               OpBranchConditional %true %24 %30
         %24 = OpLabel
               OpBranch %30
         %30 = OpLabel
        %101 = OpPhi %bool %true %24 %false %20
        %102 = OpCopyObject %bool %101
               OpBranchConditional %true %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 true  # -> $B3
          }
          $B5: {  # false
            continue false  # -> $B3
          }
        }
        unreachable
      }
      $B3 (%2:bool): {  # continuing
        %3:bool = let %2
        break_if true  # -> [t: exit_loop loop_1, f: $B2]
      }
    }
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, Phi_Loop_WithMultiblockContinue) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %bool = OpTypeBool
       %true = OpConstantTrue %bool
      %false = OpConstantFalse %bool
    %ep_type = OpTypeFunction %void
       %main = OpFunction %void None %ep_type
         %10 = OpLabel
               OpBranch %20
         %20 = OpLabel
        %101 = OpPhi %bool %true %10 %false %31
               OpLoopMerge %99 %30 None
               OpBranchConditional %101 %99 %25
         %25 = OpLabel
               OpBranch %30
         %30 = OpLabel
        %102 = OpCopyObject %bool %101
               OpBranch %31
         %31 = OpLabel
               OpBranch %20
         %99 = OpLabel
               OpReturn
               OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    loop [i: $B2, b: $B3, c: $B4] {  # loop_1
      $B2: {  # initializer
        next_iteration true  # -> $B3
      }
      $B3 (%2:bool): {  # body
        if %2 [t: $B5, f: $B6] {  # if_1
          $B5: {  # true
            exit_loop  # loop_1
          }
          $B6: {  # false
            continue  # -> $B4
          }
        }
        unreachable
      }
      $B4: {  # continuing
        %3:bool = let %2
        next_iteration false  # -> $B3
      }
    }
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, Phi_Loop_BranchConditionalBreak) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %bool = OpTypeBool
        %int = OpTypeInt 32 1
       %true = OpConstantTrue %bool
      %int_1 = OpConstant %int 1
      %int_2 = OpConstant %int 2
      %int_3 = OpConstant %int 3
      %int_4 = OpConstant %int 4
    %ep_type = OpTypeFunction %void
       %main = OpFunction %void None %ep_type
 %main_start = OpLabel
               OpBranch %1
          %1 = OpLabel
               OpLoopMerge %99 %20 None
               OpBranchConditional %true %50 %20
         %50 = OpLabel
          %6 = OpIAdd %int %int_1 %int_2
               OpBranch %99
         %20 = OpLabel
               OpBranch %1
         %99 = OpLabel
          %8 = OpPhi %int %6 %50
          %9 = OpIAdd %int %8 %8
               OpReturn
               OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    %2:i32 = loop [b: $B2, c: $B3] {  # loop_1
      $B2: {  # body
        if true [t: $B4, f: $B5] {  # if_1
          $B4: {  # true
            %3:i32 = spirv.add<i32> 1i, 2i
            exit_loop %3  # loop_1
          }
          $B5: {  # false
            continue  # -> $B3
          }
        }
        unreachable
      }
      $B3: {  # continuing
        next_iteration  # -> $B2
      }
    }
    %4:i32 = spirv.add<i32> %2, %2
    ret
  }
}
)");
}

// Phis must act as if they are simultaneously assigned. %101 and %102 should exchange values on
// each iteration, and never have the same value.
TEST_F(SpirvParserTest, Phi_SimultaneousAssignment) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
               OpName %101 "default_true"
               OpName %102 "default_false"
       %void = OpTypeVoid
       %bool = OpTypeBool
       %true = OpConstantTrue %bool
      %false = OpConstantFalse %bool
    %ep_type = OpTypeFunction %void
       %main = OpFunction %void None %ep_type
         %10 = OpLabel
               OpBranch %20
         %20 = OpLabel
        %101 = OpPhi %bool %true %10 %102 %20
        %102 = OpPhi %bool %false %10 %101 %20
               OpLoopMerge %99 %20 None
               OpBranchConditional %true %99 %20
         %99 = OpLabel
               OpReturn
               OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    loop [i: $B2, b: $B3, c: $B4] {  # loop_1
      $B2: {  # initializer
        next_iteration true, false  # -> $B3
      }
      $B3 (%2:bool, %3:bool): {  # body
        if true [t: $B5, f: $B6] {  # if_1
          $B5: {  # true
            exit_loop  # loop_1
          }
          $B6: {  # false
            continue  # -> $B4
          }
        }
        unreachable
      }
      $B4: {  # continuing
        next_iteration %3, %2  # -> $B3
      }
    }
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, Phi_SingleBlockLoopIndex) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
               OpName %2 "computed"
               OpName %3 "copied"
       %void = OpTypeVoid
       %bool = OpTypeBool
       %uint = OpTypeInt 32 0
        %pty = OpTypePointer Private %uint
          %1 = OpVariable %pty Private
    %boolpty = OpTypePointer Private %bool
          %7 = OpVariable %boolpty Private
          %8 = OpVariable %boolpty Private
     %uint_0 = OpConstant %uint 0
     %uint_1 = OpConstant %uint 1
    %ep_type = OpTypeFunction %void
       %main = OpFunction %void None %ep_type
 %main_start = OpLabel
               OpBranch %10
; Use an outer loop to show we put the new variable in the
; smallest enclosing scope.
         %10 = OpLabel
        %101 = OpLoad %bool %7
        %102 = OpLoad %bool %8
        %103 = OpIAdd %uint %uint_0 %uint_0
               OpLoopMerge %99 %89 None
               OpBranchConditional %101 %99 %20
         %20 = OpLabel
          %2 = OpPhi %uint %103 %10 %4 %20  ; gets computed value
          %3 = OpPhi %uint %uint_1 %10 %3 %20  ; gets itself
          %4 = OpIAdd %uint %2 %uint_1
               OpLoopMerge %79 %20 None
               OpBranchConditional %102 %79 %20
         %79 = OpLabel
               OpBranch %89
         %89 = OpLabel
               OpBranch %10
         %99 = OpLabel
               OpReturn
               OpFunctionEnd
)",
              R"(
$B1: {  # root
  %1:ptr<private, u32, read_write> = var undef
  %2:ptr<private, bool, read_write> = var undef
  %3:ptr<private, bool, read_write> = var undef
}

%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B2: {
    loop [b: $B3, c: $B4] {  # loop_1
      $B3: {  # body
        %5:bool = load %2
        %6:bool = load %3
        %7:u32 = spirv.add<u32> 0u, 0u
        if %5 [t: $B5, f: $B6] {  # if_1
          $B5: {  # true
            exit_loop  # loop_1
          }
          $B6: {  # false
            loop [i: $B7, b: $B8, c: $B9] {  # loop_2
              $B7: {  # initializer
                next_iteration %7, 1u  # -> $B8
              }
              $B8 (%8:u32, %9:u32): {  # body
                %10:u32 = spirv.add<u32> %8, 1u
                if %6 [t: $B10, f: $B11] {  # if_2
                  $B10: {  # true
                    exit_loop  # loop_2
                  }
                  $B11: {  # false
                    continue  # -> $B9
                  }
                }
                unreachable
              }
              $B9: {  # continuing
                next_iteration %10, %9  # -> $B8
              }
            }
            continue  # -> $B4
          }
        }
        unreachable
      }
      $B4: {  # continuing
        next_iteration  # -> $B3
      }
    }
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, Phi_MultiBlockLoopIndex) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %bool = OpTypeBool
       %uint = OpTypeInt 32 0
        %pty = OpTypePointer Private %uint
          %1 = OpVariable %pty Private
    %boolpty = OpTypePointer Private %bool
          %7 = OpVariable %boolpty Private
          %8 = OpVariable %boolpty Private
     %uint_0 = OpConstant %uint 0
     %uint_1 = OpConstant %uint 1
    %ep_type = OpTypeFunction %void
       %main = OpFunction %void None %ep_type
          %5 = OpLabel
               OpBranch %10
; Use an outer loop to show we put the new variable in the
; smallest enclosing scope.
         %10 = OpLabel
        %101 = OpLoad %bool %7
        %102 = OpLoad %bool %8
               OpLoopMerge %99 %89 None
               OpBranchConditional %101 %99 %20
         %20 = OpLabel
          %2 = OpPhi %uint %uint_0 %10 %4 %30  ; gets computed value
          %3 = OpPhi %uint %uint_1 %10 %3 %30  ; gets itself
               OpLoopMerge %79 %30 None
               OpBranchConditional %102 %79 %30
         %30 = OpLabel  ; continue target for inner loop
          %4 = OpIAdd %uint %2 %uint_1
               OpBranch %20
         %79 = OpLabel  ; merge for inner loop
               OpBranch %89
         %89 = OpLabel  ; continue target for outer loop
               OpBranch %10
         %99 = OpLabel  ; merge for outer loop
               OpReturn
               OpFunctionEnd
)",
              R"(
$B1: {  # root
  %1:ptr<private, u32, read_write> = var undef
  %2:ptr<private, bool, read_write> = var undef
  %3:ptr<private, bool, read_write> = var undef
}

%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B2: {
    loop [b: $B3, c: $B4] {  # loop_1
      $B3: {  # body
        %5:bool = load %2
        %6:bool = load %3
        if %5 [t: $B5, f: $B6] {  # if_1
          $B5: {  # true
            exit_loop  # loop_1
          }
          $B6: {  # false
            loop [i: $B7, b: $B8, c: $B9] {  # loop_2
              $B7: {  # initializer
                next_iteration 0u, 1u  # -> $B8
              }
              $B8 (%7:u32, %8:u32): {  # body
                if %6 [t: $B10, f: $B11] {  # if_2
                  $B10: {  # true
                    exit_loop  # loop_2
                  }
                  $B11: {  # false
                    continue  # -> $B9
                  }
                }
                unreachable
              }
              $B9: {  # continuing
                %9:u32 = spirv.add<u32> %7, 1u
                next_iteration %9, %8  # -> $B8
              }
            }
            continue  # -> $B4
          }
        }
        unreachable
      }
      $B4: {  # continuing
        next_iteration  # -> $B3
      }
    }
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, Phi_ValueFromLoopBodyAndContinuing) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %bool = OpTypeBool
       %uint = OpTypeInt 32 0
        %pty = OpTypePointer Private %uint
          %1 = OpVariable %pty Private
    %boolpty = OpTypePointer Private %bool
         %17 = OpVariable %boolpty Private
     %uint_0 = OpConstant %uint 0
     %uint_1 = OpConstant %uint 1
       %true = OpConstantTrue %bool
    %ep_type = OpTypeFunction %void
       %main = OpFunction %void None %ep_type
          %9 = OpLabel
        %101 = OpLoad %bool %17
               OpBranch %10
; Use an outer loop to show we put the new variable in the
; smallest enclosing scope.
         %10 = OpLabel
               OpLoopMerge %99 %89 None
               OpBranch %20
         %20 = OpLabel
          %2 = OpPhi %uint %uint_0 %10 %4 %30  ; gets computed value
          %5 = OpPhi %uint %uint_1 %10 %7 %30
          %4 = OpIAdd %uint %2 %uint_1 ; define %4
          %6 = OpIAdd %uint %4 %uint_1 ; use %4
               OpLoopMerge %79 %30 None
               OpBranchConditional %101 %79 %30
         %30 = OpLabel
          %7 = OpIAdd %uint %4 %6 ; use %4 again
          %8 = OpCopyObject %uint %5 ; use %5
               OpBranchConditional %true %20 %79
         %79 = OpLabel
               OpBranch %89
         %89 = OpLabel
               OpBranchConditional %true %10 %99
         %99 = OpLabel
               OpReturn
               OpFunctionEnd
)",
              R"(
$B1: {  # root
  %1:ptr<private, u32, read_write> = var undef
  %2:ptr<private, bool, read_write> = var undef
}

%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B2: {
    %4:bool = load %2
    loop [b: $B3, c: $B4] {  # loop_1
      $B3: {  # body
        loop [i: $B5, b: $B6, c: $B7] {  # loop_2
          $B5: {  # initializer
            next_iteration 0u, 1u  # -> $B6
          }
          $B6 (%5:u32, %6:u32): {  # body
            %7:u32 = spirv.add<u32> %5, 1u
            %8:u32 = spirv.add<u32> %7, 1u
            if %4 [t: $B8, f: $B9] {  # if_1
              $B8: {  # true
                exit_loop  # loop_2
              }
              $B9: {  # false
                continue  # -> $B7
              }
            }
            unreachable
          }
          $B7: {  # continuing
            %9:u32 = spirv.add<u32> %7, %8
            %10:u32 = let %6
            %11:bool = not true
            break_if %11 next_iteration: [ %7, %9 ]  # -> [t: exit_loop loop_2, f: $B6]
          }
        }
        continue  # -> $B4
      }
      $B4: {  # continuing
        %12:bool = not true
        break_if %12  # -> [t: exit_loop loop_1, f: $B3]
      }
    }
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, Phi_FromElseAndThen) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %bool = OpTypeBool
       %uint = OpTypeInt 32 0
        %pty = OpTypePointer Private %uint
          %1 = OpVariable %pty Private
    %boolpty = OpTypePointer Private %bool
          %7 = OpVariable %boolpty Private
          %8 = OpVariable %boolpty Private
     %uint_0 = OpConstant %uint 0
     %uint_1 = OpConstant %uint 1
     %uint_3 = OpConstant %uint 3
    %ep_type = OpTypeFunction %void
       %main = OpFunction %void None %ep_type
          %5 = OpLabel
        %101 = OpLoad %bool %7
        %102 = OpLoad %bool %8
               OpBranch %10
; Use an outer loop to show we put the new variable in the
; smallest enclosing scope.
         %10 = OpLabel
               OpLoopMerge %99 %89 None
               OpBranchConditional %101 %99 %20
         %20 = OpLabel ; if seleciton
               OpSelectionMerge %79 None
               OpBranchConditional %102 %30 %40
         %30 = OpLabel
               OpBranch %89
         %40 = OpLabel
               OpBranch %89
         %79 = OpLabel ; disconnected selection merge node
               OpBranch %89
         %89 = OpLabel
          %2 = OpPhi %uint %uint_0 %30 %uint_1 %40 %uint_3 %79
               OpStore %1 %2
               OpBranch %10
         %99 = OpLabel
               OpReturn
               OpFunctionEnd
)",
              R"(
$B1: {  # root
  %1:ptr<private, u32, read_write> = var undef
  %2:ptr<private, bool, read_write> = var undef
  %3:ptr<private, bool, read_write> = var undef
}

%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B2: {
    %5:bool = load %2
    %6:bool = load %3
    loop [b: $B3, c: $B4] {  # loop_1
      $B3: {  # body
        if %5 [t: $B5, f: $B6] {  # if_1
          $B5: {  # true
            exit_loop  # loop_1
          }
          $B6: {  # false
            if %6 [t: $B7, f: $B8] {  # if_2
              $B7: {  # true
                continue 0u  # -> $B4
              }
              $B8: {  # false
                continue 1u  # -> $B4
              }
            }
            continue 3u  # -> $B4
          }
        }
        unreachable
      }
      $B4 (%7:u32): {  # continuing
        store %1, %7
        next_iteration  # -> $B3
      }
    }
    ret
  }
}
)");
}

TEST_F(SpirvParserTest, Phi_FromHeaderAndThen) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %bool = OpTypeBool
       %uint = OpTypeInt 32 0
        %pty = OpTypePointer Private %uint
          %1 = OpVariable %pty Private
    %boolpty = OpTypePointer Private %bool
          %7 = OpVariable %boolpty Private
          %8 = OpVariable %boolpty Private
     %uint_0 = OpConstant %uint 0
     %uint_1 = OpConstant %uint 1
    %ep_type = OpTypeFunction %void
       %main = OpFunction %void None %ep_type
          %5 = OpLabel
        %101 = OpLoad %bool %7
        %102 = OpLoad %bool %8
               OpBranch %10
; Use an outer loop to show we put the new variable in the
; smallest enclosing scope.
         %10 = OpLabel
               OpLoopMerge %99 %89 None
               OpBranchConditional %101 %99 %20
         %20 = OpLabel ; if seleciton
               OpSelectionMerge %79 None
               OpBranchConditional %102 %30 %89
         %30 = OpLabel
               OpBranch %89
         %79 = OpLabel ; disconnected selection merge node
               OpUnreachable
         %89 = OpLabel
          %2 = OpPhi %uint %uint_0 %20 %uint_1 %30
          %3 = OpPhi %uint %uint_1 %20 %uint_0 %30
          %4 = OpIAdd %uint %2 %3
               OpStore %1 %4
               OpBranch %10
         %99 = OpLabel
               OpReturn
               OpFunctionEnd
)",
              R"(
$B1: {  # root
  %1:ptr<private, u32, read_write> = var undef
  %2:ptr<private, bool, read_write> = var undef
  %3:ptr<private, bool, read_write> = var undef
}

%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B2: {
    %5:bool = load %2
    %6:bool = load %3
    loop [b: $B3, c: $B4] {  # loop_1
      $B3: {  # body
        if %5 [t: $B5, f: $B6] {  # if_1
          $B5: {  # true
            exit_loop  # loop_1
          }
          $B6: {  # false
            if %6 [t: $B7, f: $B8] {  # if_2
              $B7: {  # true
                continue 1u, 0u  # -> $B4
              }
              $B8: {  # false
                continue 0u, 1u  # -> $B4
              }
            }
            unreachable
          }
        }
        unreachable
      }
      $B4 (%7:u32, %8:u32): {  # continuing
        %9:u32 = spirv.add<u32> %7, %8
        store %1, %9
        next_iteration  # -> $B3
      }
    }
    ret
  }
}
)");
}

// If the only use of a combinatorially computed ID is as the value in an OpPhi, then we still have
// to emit it.  The algorithm fix is to always count uses in Phis. This is the reduced case from the
// bug report.
//
// * The only use of %12 is in the phi.
// * The only use of %11 is in %12.
// * Both definitions need to be emitted to the output.
//
// https://crbug.com/215
TEST_F(SpirvParserTest, Phi_UseInPhiCountsAsUse) {
    EXPECT_IR(
        R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %bool = OpTypeBool
       %true = OpConstantTrue %bool
    %ep_type = OpTypeFunction %void
       %main = OpFunction %void None %ep_type
         %10 = OpLabel
         %11 = OpLogicalAnd %bool %true %true
         %12 = OpLogicalNot %bool %11  ;
               OpSelectionMerge %99 None
               OpBranchConditional %true %20 %99
         %20 = OpLabel
               OpBranch %99
         %99 = OpLabel
        %101 = OpPhi %bool %11 %10 %12 %20
        %102 = OpCopyObject %bool %101  ;; ensure a use of %101
               OpReturn
               OpFunctionEnd
)",
        R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    %2:bool = and true, true
    %3:bool = not %2
    %4:bool = if true [t: $B2, f: $B3] {  # if_1
      $B2: {  # true
        exit_if %3  # if_1
      }
      $B3: {  # false
        exit_if %2  # if_1
      }
    }
    %5:bool = let %4
    ret
  }
}
)");
}

// Value %999 is defined deep in control flow, then we arrange for it to dominate the backedge of
// the outer loop. The %999 value is then fed back into the phi in the loop header.  So %999 needs
// to be hoisted out of the loop.  The phi assignment needs to use the hoisted variable. The hoisted
// variable needs to be placed such that its scope encloses that phi in the header of the outer
// loop. The compiler needs to "see" that there is an implicit use of %999 in the backedge block of
// that outer loop.
//
// https://crbug.com/1649
TEST_F(SpirvParserTest, Phi_PhiInLoopHeader_FedByHoistedVar_PhiUnused) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %bool = OpTypeBool
       %true = OpConstantTrue %bool
    %ep_type = OpTypeFunction %void
       %main = OpFunction %void None %ep_type
         %10 = OpLabel
               OpBranch %20
         %20 = OpLabel
        %101 = OpPhi %bool %true %10 %999 %80
               OpLoopMerge %99 %80 None
               OpBranchConditional %true %30 %99
         %30 = OpLabel
               OpSelectionMerge %50 None
               OpBranchConditional %true %40 %50
         %40 = OpLabel
        %999 = OpCopyObject %bool %true
               OpBranch %60
         %50 = OpLabel
               OpReturn
         %60 = OpLabel ; if merge
               OpBranch %80
         %80 = OpLabel ; continue target
               OpBranch %20
         %99 = OpLabel
               OpReturn
               OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    loop [i: $B2, b: $B3, c: $B4] {  # loop_1
      $B2: {  # initializer
        next_iteration true  # -> $B3
      }
      $B3 (%2:bool): {  # body
        if true [t: $B5, f: $B6] {  # if_1
          $B5: {  # true
            if true [t: $B7, f: $B8] {  # if_2
              $B7: {  # true
                %3:bool = let true
                continue %3  # -> $B4
              }
              $B8: {  # false
                exit_if  # if_2
              }
            }
            ret
          }
          $B6: {  # false
            exit_loop  # loop_1
          }
        }
        unreachable
      }
      $B4 (%4:bool): {  # continuing
        next_iteration %4  # -> $B3
      }
    }
    ret
  }
}
)");
}

// Value %999 is defined deep in control flow, then we arrange for it to dominate the backedge of
// the outer loop. The %999 value is then fed back into the phi in the loop header.  So %999 needs
// to be hoisted out of the loop.  The phi assignment needs to use the hoisted variable. The hoisted
// variable needs to be placed such that its scope encloses that phi in the header of the outer
// loop. The compiler needs to "see" that there is an implicit use of %999 in the backedge block of
// that outer loop.
//
// https://crbug.com/1649
TEST_F(SpirvParserTest, Phi_PhiInLoopHeader_FedByHoistedVar_PhiUsed) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %bool = OpTypeBool
       %true = OpConstantTrue %bool
    %ep_type = OpTypeFunction %void
       %main = OpFunction %void None %ep_type
         %10 = OpLabel
               OpBranch %20
         %20 = OpLabel
        %101 = OpPhi %bool %true %10 %999 %80
               OpLoopMerge %99 %80 None
               OpBranchConditional %true %30 %99
         %30 = OpLabel
               OpSelectionMerge %50 None
               OpBranchConditional %true %40 %50
         %40 = OpLabel
        %999 = OpCopyObject %bool %true
               OpBranch %60
         %50 = OpLabel
               OpReturn
         %60 = OpLabel ; if merge
               OpBranch %80
         %80 = OpLabel ; continue target
               OpBranch %20
         %99 = OpLabel
       %1000 = OpCopyObject %bool %101
               OpReturn
               OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    %2:bool = loop [i: $B2, b: $B3, c: $B4] {  # loop_1
      $B2: {  # initializer
        next_iteration true  # -> $B3
      }
      $B3 (%3:bool): {  # body
        if true [t: $B5, f: $B6] {  # if_1
          $B5: {  # true
            if true [t: $B7, f: $B8] {  # if_2
              $B7: {  # true
                %4:bool = let true
                continue %4  # -> $B4
              }
              $B8: {  # false
                exit_if  # if_2
              }
            }
            ret
          }
          $B6: {  # false
            exit_loop %3  # loop_1
          }
        }
        unreachable
      }
      $B4 (%5:bool): {  # continuing
        next_iteration %5  # -> $B3
      }
    }
    %6:bool = let %2
    ret
  }
}
)");
}

// This is a reduction of one of the hard parts of test case
// vk-gl-cts/graphicsfuzz/stable-binarysearch-tree-false-if-discard-loop/1.spvasm
// In particular, see the data flow around %114 in that case.
//
// Here value %999 is is a *phi* defined deep in control flow, then we arrange for it to dominate
// the backedge of the outer loop. The %999 value is then fed back into the phi in the loop header.
// The variable generated to hold the %999 value needs to be placed such that its scope encloses
// that phi in the header of the outer loop. The compiler needs to "see" that there is an implicit
// use of %999 in the backedge block of that outer loop.
//
// https://crbug.com/1649
TEST_F(SpirvParserTest, Phi_PhiInLoopHeader_FedByPhi_PhiUnused) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %bool = OpTypeBool
       %true = OpConstantTrue %bool
      %false = OpConstantFalse %bool
    %ep_type = OpTypeFunction %void
       %main = OpFunction %void None %ep_type
         %10 = OpLabel
               OpBranch %20
         %20 = OpLabel
        %101 = OpPhi %bool %true %10 %999 %80
               OpLoopMerge %99 %80 None
               OpBranchConditional %true %99 %30
         %30 = OpLabel
               OpLoopMerge %70 %60 None
               OpBranch %40
         %40 = OpLabel
               OpBranchConditional %true %60 %50
         %50 = OpLabel
               OpBranch %60
         %60 = OpLabel ; inner continue
        %999 = OpPhi %bool %true %40 %false %50
               OpBranchConditional %true %70 %30
         %70 = OpLabel  ; inner merge
               OpBranch %80
         %80 = OpLabel ; outer continue target
               OpBranch %20
         %99 = OpLabel
               OpReturn
               OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    loop [i: $B2, b: $B3, c: $B4] {  # loop_1
      $B2: {  # initializer
        next_iteration true  # -> $B3
      }
      $B3 (%2:bool): {  # body
        if true [t: $B5, f: $B6] {  # if_1
          $B5: {  # true
            exit_loop  # loop_1
          }
          $B6: {  # false
            %3:bool = loop [b: $B7, c: $B8] {  # loop_2
              $B7: {  # body
                if true [t: $B9, f: $B10] {  # if_2
                  $B9: {  # true
                    continue true  # -> $B8
                  }
                  $B10: {  # false
                    continue false  # -> $B8
                  }
                }
                unreachable
              }
              $B8 (%4:bool): {  # continuing
                break_if true exit_loop: [ %4 ]  # -> [t: exit_loop loop_2, f: $B7]
              }
            }
            continue %3  # -> $B4
          }
        }
        unreachable
      }
      $B4 (%5:bool): {  # continuing
        next_iteration %5  # -> $B3
      }
    }
    ret
  }
}
)");
}

// This is a reduction of one of the hard parts of test case
// vk-gl-cts/graphicsfuzz/stable-binarysearch-tree-false-if-discard-loop/1.spvasm
// In particular, see the data flow around %114 in that case.
//
// Here value %999 is is a *phi* defined deep in control flow, then we arrange for it to dominate
// the backedge of the outer loop. The %999 value is then fed back into the phi in the loop header.
// The variable generated to hold the %999 value needs to be placed such that its scope encloses
// that phi in the header of the outer loop. The compiler needs to "see" that there is an implicit
// use of %999 in the backedge block of that outer loop.
//
// https://crbug.com/1649
TEST_F(SpirvParserTest, Phi_PhiInLoopHeader_FedByPhi_PhiUsed) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %bool = OpTypeBool
       %true = OpConstantTrue %bool
      %false = OpConstantFalse %bool
    %ep_type = OpTypeFunction %void
       %main = OpFunction %void None %ep_type
         %10 = OpLabel
               OpBranch %20
         %20 = OpLabel
        %101 = OpPhi %bool %true %10 %999 %80
               OpLoopMerge %99 %80 None
               OpBranchConditional %true %99 %30
         %30 = OpLabel
               OpLoopMerge %70 %60 None
               OpBranch %40
         %40 = OpLabel
               OpBranchConditional %true %60 %50
         %50 = OpLabel
               OpBranch %60
         %60 = OpLabel ; inner continue
        %999 = OpPhi %bool %true %40 %false %50
               OpBranchConditional %true %70 %30
         %70 = OpLabel  ; inner merge
               OpBranch %80
         %80 = OpLabel ; outer continue target
               OpBranch %20
         %99 = OpLabel
       %1000 = OpCopyObject %bool %101
               OpReturn
               OpFunctionEnd
)",
              R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
  $B1: {
    %2:bool = loop [i: $B2, b: $B3, c: $B4] {  # loop_1
      $B2: {  # initializer
        next_iteration true  # -> $B3
      }
      $B3 (%3:bool): {  # body
        if true [t: $B5, f: $B6] {  # if_1
          $B5: {  # true
            exit_loop %3  # loop_1
          }
          $B6: {  # false
            %4:bool = loop [b: $B7, c: $B8] {  # loop_2
              $B7: {  # body
                if true [t: $B9, f: $B10] {  # if_2
                  $B9: {  # true
                    continue true  # -> $B8
                  }
                  $B10: {  # false
                    continue false  # -> $B8
                  }
                }
                unreachable
              }
              $B8 (%5:bool): {  # continuing
                break_if true exit_loop: [ %5 ]  # -> [t: exit_loop loop_2, f: $B7]
              }
            }
            continue %4  # -> $B4
          }
        }
        unreachable
      }
      $B4 (%6:bool): {  # continuing
        next_iteration %6  # -> $B3
      }
    }
    %7:bool = let %2
    ret
  }
}
)");
}

// A phi in an unreachable block may have no operands.
TEST_F(SpirvParserTest, Phi_UnreachableLoopMerge) {
    EXPECT_IR(R"(
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 1 1 1
       %void = OpTypeVoid
       %uint = OpTypeInt 32 0
    %ep_type = OpTypeFunction %void
       %main = OpFunction %void None %ep_type
         %10 = OpLabel
               OpBranch %99
         %99 = OpLabel
               OpLoopMerge %101 %99 None
               OpBranch %99
        %101 = OpLabel
        %102 = OpPhi %uint
               OpUnreachable
               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
      }
    }
    unreachable
  }
}
)");
}

}  // namespace
}  // namespace tint::spirv::reader