Addition of OpSwitchBranches.ll, OpSwitchUnreachable.ll, Two_OpSwitch_same_register.ll and switch-range-check.ll

test/branching/OpSwitchBranches.ll

@@ -0,0 +1,49 @@
+; RUN: llvm-as < %s -o %t.bc
+; RUN: llvm-spirv %t.bc -o %t.spv
+; RUN: spirv-val %t.spv
+; RUN: llvm-spirv -to-text %t.spv -o - | FileCheck %s
+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
+target triple = "spir64-unknown-unknown"
+
+define i32 @test_switch_branches(i32 %a) {
+entry:
+  %alloc = alloca i32
+; CHECK:      Switch [[#]] [[#DEFAULT:]] 1 [[#CASE1:]] 2 [[#CASE2:]] 3 [[#CASE3:]]
+  switch i32 %a, label %default [
+    i32 1, label %case1
+    i32 2, label %case2
+    i32 3, label %case3
+  ]
+
+case1:
+  store i32 1, ptr %alloc
+  br label %end
+
+case2:
+  store i32 2, ptr %alloc
+  br label %end
+
+case3:
+  store i32 3, ptr %alloc
+  br label %end
+
+default:
+  store i32 0, ptr %alloc
+  br label %end
+
+end:
+  %result = load i32, ptr %alloc
+  ret i32 %result
+
+; CHECK:      Label [[#CASE1]] 
+; CHECK:      Branch [[#END:]]
+; CHECK:      Label [[#CASE2]] 
+; CHECK:      Branch [[#END]]
+; CHECK:      Label [[#CASE3]]
+; CHECK:      Branch [[#END:]]
+; CHECK:      Label [[#DEFAULT]] 
+; CHECK:      Branch [[#END]]
+; CHECK:      Label [[#END]] 
+; CHECK:      ReturnValue
+}

test/branching/OpSwitchChar.ll

@@ -0,0 +1,57 @@
+;; __kernel void test_switch(__global int* res, uchar val)
+;; {
+;;   switch(val)
+;;   {
+;;   case 0:
+;;     *res = 1;
+;;     break;
+;;   case 1:
+;;     *res = 2;
+;;     break;
+;;   case 2:
+;;     *res = 3;
+;;     break;
+;;   }
+;; }
+
+; RUN: llvm-as < %s -o %t.bc
+; RUN: llvm-spirv %t.bc -o %t.spv
+; RUN: spirv-val %t.spv
+; RUN: llvm-spirv -to-text %t.spv -o - | FileCheck %s
+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
+target triple = "spir64-unknown-unknown"
+
+; CHECK: Switch [[#]] [[#]] 0 [[#]] 1 [[#]] 2 [[#]]
+
+define spir_kernel void @test_switch(i32 addrspace(1)* %res, i8 zeroext %val) {
+entry:
+  %res.addr = alloca i32 addrspace(1)*, align 4
+  %val.addr = alloca i8, align 1
+  store i32 addrspace(1)* %res, i32 addrspace(1)** %res.addr, align 4
+  store i8 %val, i8* %val.addr, align 1
+  %0 = load i8, i8* %val.addr, align 1
+  switch i8 %0, label %sw.epilog [
+    i8 0, label %sw.bb
+    i8 1, label %sw.bb1
+    i8 2, label %sw.bb2
+  ]
+
+sw.bb:                                            ; preds = %entry
+  %1 = load i32 addrspace(1)*, i32 addrspace(1)** %res.addr, align 4
+  store i32 1, i32 addrspace(1)* %1, align 4
+  br label %sw.epilog
+
+sw.bb1:                                           ; preds = %entry
+  %2 = load i32 addrspace(1)*, i32 addrspace(1)** %res.addr, align 4
+  store i32 2, i32 addrspace(1)* %2, align 4
+  br label %sw.epilog
+
+sw.bb2:                                           ; preds = %entry
+  %3 = load i32 addrspace(1)*, i32 addrspace(1)** %res.addr, align 4
+  store i32 3, i32 addrspace(1)* %3, align 4
+  br label %sw.epilog
+
+sw.epilog:                                        ; preds = %entry, %sw.bb2, %sw.bb1, %sw.bb
+  ret void
+}

test/branching/OpSwitchUnreachable.ll

@@ -0,0 +1,26 @@
+; RUN: llvm-as < %s -o %t.bc
+; RUN: llvm-spirv %t.bc -o %t.spv
+; RUN: spirv-val %t.spv
+; RUN: llvm-spirv -to-text %t.spv -o - | FileCheck %s
+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
+target triple = "spir64-unknown-unknown"
+
+define void @test_switch_with_unreachable_block(i1 %a) {
+  %value = zext i1 %a to i32
+; CHECK:      Switch [[#]] [[#UNREACHABLE:]] 0 [[#REACHABLE:]] 1 [[#REACHABLE:]]
+  switch i32 %value, label %unreachable [
+    i32 0, label %reachable
+    i32 1, label %reachable
+  ]
+
+; CHECK: Label [[#REACHABLE]]
+reachable:
+; CHECK: Return
+  ret void
+
+; CHECK: Label [[#UNREACHABLE]]
+; CHECK: Unreachable
+unreachable:
+  unreachable
+}

test/branching/Two_OpSwitch_same_register.ll

@@ -0,0 +1,54 @@
+; RUN: llvm-as < %s -o %t.bc
+; RUN: llvm-spirv %t.bc -o %t.spv
+; RUN: spirv-val %t.spv
+; RUN: llvm-spirv -to-text %t.spv -o - | FileCheck %s
+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
+target triple = "spir64-unknown-unknown"
+
+define spir_kernel void @test_two_switch_same_register(i32 %value) {
+; CHECK: Switch [[#REGISTER:]] [[#DEFAULT1:]] 1 [[#CASE1:]] 0 [[#CASE2:]]
+  switch i32 %value, label %default1 [
+    i32 1, label %case1
+    i32 0, label %case2
+  ]
+
+case1:
+  br label %default1
+
+case2:
+  br label %default1
+
+default1:
+  switch i32 %value, label %default2 [
+    i32 0, label %case3
+    i32 1, label %case4
+  ]
+
+case3:
+  br label %default2
+
+case4:
+  br label %default2
+
+default2:
+  ret void
+
+; CHECK:      Label [[#CASE1]]
+; CHECK-NEXT: Branch [[#DEFAULT1]]
+
+; CHECK:      Label [[#CASE2]]
+; CHECK-NEXT: Branch [[#DEFAULT1]]
+
+; CHECK:      Label [[#DEFAULT1]]
+; CHECK-NEXT: Switch [[#REGISTER]] [[#DEFAULT2:]] 0 [[#CASE3:]] 1 [[#CASE4:]]
+
+; CHECK:      Label [[#CASE3]]
+; CHECK-NEXT: Branch [[#DEFAULT2]]
+
+; CHECK:      Label [[#CASE4]]
+; CHECK-NEXT: Branch [[#DEFAULT2]]
+
+; CHECK:      Label [[#DEFAULT2]]
+; CHECK-NEXT: Return
+}

test/branching/switch-range-check.ll

@@ -0,0 +1,87 @@
+; RUN: llvm-as < %s -o %t.bc
+; RUN: llvm-spirv %t.bc -o %t.spv
+; RUN: spirv-val %t.spv
+; RUN: llvm-spirv -to-text %t.spv -o - | FileCheck %s
+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
+target triple = "spir64-unknown-unknown"
+
+; CHECK: BranchConditional [[#]] [[#if_then:]] [[#if_end:]]
+; CHECK: Label [[#if_then]]
+; CHECK: Branch [[#if_end]]
+; CHECK: Label [[#if_end]]
+; CHECK: Phi [[#]] [[#Var:]]
+; CHECK: 27 Switch [[#Var]]
+; CHECK-COUNT-10: Label
+; CHECK: Label [[#epilog:]]
+; CHECK: Branch [[#exit:]]
+; CHECK: Label [[#exit]] 
+; CHECK: Return
+; CHECK-NOT: Label
+; CHECK: FunctionEnd
+
+define spir_func void @foo(i64 noundef %addr, i64 noundef %as) {
+entry:
+  %src = inttoptr i64 %as to ptr addrspace(4)
+  %val = load i8, ptr addrspace(4) %src
+  %cmp = icmp sgt i8 %val, 0
+  br i1 %cmp, label %if.then, label %if.end
+
+if.then:
+  %add.ptr = getelementptr inbounds i8, ptr addrspace(4) %src, i64 1
+  %cond = load i8, ptr addrspace(4) %add.ptr
+  br label %if.end
+
+if.end:
+  %swval = phi i8 [ %cond, %if.then ], [ %val, %entry ]
+  switch i8 %swval, label %sw.default [
+    i8 -127, label %sw.epilog
+    i8 -126, label %sw.bb3
+    i8 -125, label %sw.bb4
+    i8 -111, label %sw.bb5
+    i8 -110, label %sw.bb6
+    i8 -109, label %sw.bb7
+    i8 -15, label %sw.bb8
+    i8 -14, label %sw.bb8
+    i8 -13, label %sw.bb8
+    i8 -124, label %sw.bb9
+    i8 -95, label %sw.bb10
+    i8 -123, label %sw.bb11
+  ]
+
+sw.bb3:
+  br label %sw.epilog
+
+sw.bb4:
+  br label %sw.epilog
+
+sw.bb5:
+  br label %sw.epilog
+
+sw.bb6:
+  br label %sw.epilog
+
+sw.bb7:
+  br label %sw.epilog
+
+sw.bb8:
+  br label %sw.epilog
+
+sw.bb9:
+  br label %sw.epilog
+
+sw.bb10:
+  br label %sw.epilog
+
+sw.bb11:
+  br label %sw.epilog
+
+sw.default:
+  br label %sw.epilog
+
+sw.epilog:
+  br label %exit
+
+exit:
+  ret void
+}