diff --git a/llvm/test/CodeGen/SPIRV/LinkOnceODR.ll b/llvm/test/CodeGen/SPIRV/LinkOnceODR.ll
new file mode 100644
index 0000000000000..3fb49ac241c6e
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/LinkOnceODR.ll
@@ -0,0 +1,24 @@
+;; No extension -> no LinkOnceODR
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV-NOT: OpExtension "SPV_KHR_linkonce_odr"
+; CHECK-SPIRV-NOT: OpDecorate %[[#]] LinkageAttributes "GV" LinkOnceODR 
+; CHECK-SPIRV-NOT: OpDecorate %[[#]] LinkageAttributes "square" LinkOnceODR 
+
+@GV = linkonce_odr addrspace(1) global [3 x i32] zeroinitializer, align 4
+
+define spir_kernel void @k() {
+entry:
+  %call = call spir_func i32 @square(i32 2)
+  ret void
+}
+
+define linkonce_odr dso_local spir_func i32 @square(i32 %in) {
+entry:
+  %in.addr = alloca i32, align 4
+  store i32 %in, i32* %in.addr, align 4
+  %0 = load i32, i32* %in.addr, align 4
+  %1 = load i32, i32* %in.addr, align 4
+  %mul = mul nsw i32 %0, %1
+  ret i32 %mul
+}
diff --git a/llvm/test/CodeGen/SPIRV/OpVectorInsertDynamic.ll b/llvm/test/CodeGen/SPIRV/OpVectorInsertDynamic.ll
new file mode 100644
index 0000000000000..8da39a1315249
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/OpVectorInsertDynamic.ll
@@ -0,0 +1,16 @@
+;; uint8 foo(uint8 c, unsigned i) {
+;;   c[i] = 42;
+;;   return c;
+;; }
+
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
+
+; CHECK: %[[#TypeInt:]] = OpTypeInt 32
+; CHECK: %[[#TypeVector:]] = OpTypeVector %[[#TypeInt]] 8
+; CHECK: %[[#]] = OpVectorInsertDynamic %[[#TypeVector]]
+
+define spir_func <8 x i32> @foo(<8 x i32> %c, i32 %i) local_unnamed_addr {
+entry:
+  %vecins = insertelement <8 x i32> %c, i32 42, i32 %i
+  ret <8 x i32> %vecins
+}
diff --git a/llvm/test/CodeGen/SPIRV/atomicrmw.ll b/llvm/test/CodeGen/SPIRV/atomicrmw.ll
new file mode 100644
index 0000000000000..401b21224e7dd
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/atomicrmw.ll
@@ -0,0 +1,57 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+
+; CHECK:     %[[#Int:]] = OpTypeInt 32 0
+; CHECK-DAG: %[[#Scope_Device:]] = OpConstant %[[#Int]] 1 {{$}}
+; CHECK-DAG: %[[#MemSem_Relaxed:]] = OpConstant %[[#Int]] 0
+; CHECK-DAG: %[[#MemSem_Acquire:]] = OpConstant %[[#Int]] 2
+; CHECK-DAG: %[[#MemSem_Release:]] = OpConstant %[[#Int]] 4 {{$}}
+; CHECK-DAG: %[[#MemSem_AcquireRelease:]] = OpConstant %[[#Int]] 8
+; CHECK-DAG: %[[#MemSem_SequentiallyConsistent:]] = OpConstant %[[#Int]] 16
+; CHECK-DAG: %[[#Value:]] = OpConstant %[[#Int]] 42
+; CHECK-DAG: %[[#Float:]] = OpTypeFloat 32
+; CHECK-DAG: %[[#PointerType:]] = OpTypePointer CrossWorkgroup %[[#Int]]
+; CHECK-DAG: %[[#FPPointerType:]] = OpTypePointer CrossWorkgroup %[[#Float]]
+; CHECK-DAG: %[[#Pointer:]] = OpVariable %[[#PointerType]] CrossWorkgroup
+; CHECK-DAG: %[[#FPPointer:]] = OpVariable %[[#FPPointerType]] CrossWorkgroup
+; CHECK-DAG: %[[#FPValue:]] = OpConstant %[[#Float]] 1109917696
+
+@ui = common dso_local addrspace(1) global i32 0, align 4
+@f = common dso_local local_unnamed_addr addrspace(1) global float 0.000000e+00, align 4
+
+define dso_local spir_func void @test_atomicrmw() local_unnamed_addr {
+entry:
+  %0 = atomicrmw xchg i32 addrspace(1)* @ui, i32 42 acq_rel
+; CHECK: %[[#]] = OpAtomicExchange %[[#Int]] %[[#Pointer]] %[[#Scope_Device]] %[[#MemSem_AcquireRelease]] %[[#Value]]
+
+  %1 = atomicrmw xchg float addrspace(1)* @f, float 42.000000e+00 seq_cst
+; CHECK: %[[#]] = OpAtomicExchange %[[#Float]] %[[#FPPointer]] %[[#Scope_Device]] %[[#MemSem_SequentiallyConsistent]] %[[#FPValue]]
+
+  %2 = atomicrmw add i32 addrspace(1)* @ui, i32 42 monotonic
+; CHECK: %[[#]] = OpAtomicIAdd %[[#Int]] %[[#Pointer]] %[[#Scope_Device]] %[[#MemSem_Relaxed]] %[[#Value]]
+
+  %3 = atomicrmw sub i32 addrspace(1)* @ui, i32 42 acquire
+; CHECK: %[[#]] = OpAtomicISub %[[#Int]] %[[#Pointer]] %[[#Scope_Device]] %[[#MemSem_Acquire]] %[[#Value]]
+
+  %4 = atomicrmw or i32 addrspace(1)* @ui, i32 42 release
+; CHECK: %[[#]] = OpAtomicOr %[[#Int]] %[[#Pointer]] %[[#Scope_Device]] %[[#MemSem_Release]] %[[#Value]]
+
+  %5 = atomicrmw xor i32 addrspace(1)* @ui, i32 42 acq_rel
+; CHECK: %[[#]] = OpAtomicXor %[[#Int]] %[[#Pointer]] %[[#Scope_Device]] %[[#MemSem_AcquireRelease]] %[[#Value]]
+
+  %6 = atomicrmw and i32 addrspace(1)* @ui, i32 42 seq_cst
+; CHECK: %[[#]] = OpAtomicAnd %[[#Int]] %[[#Pointer]] %[[#Scope_Device]] %[[#MemSem_SequentiallyConsistent]] %[[#Value]]
+
+  %7 = atomicrmw max i32 addrspace(1)* @ui, i32 42 monotonic
+; CHECK: %[[#]] = OpAtomicSMax %[[#Int]] %[[#Pointer]] %[[#Scope_Device]] %[[#MemSem_Relaxed]] %[[#Value]]
+
+  %8 = atomicrmw min i32 addrspace(1)* @ui, i32 42 acquire
+; CHECK: %[[#]] = OpAtomicSMin %[[#Int]] %[[#Pointer]] %[[#Scope_Device]] %[[#MemSem_Acquire]] %[[#Value]]
+
+  %9 = atomicrmw umax i32 addrspace(1)* @ui, i32 42 release
+; CHECK: %[[#]] = OpAtomicUMax %[[#Int]] %[[#Pointer]] %[[#Scope_Device]] %[[#MemSem_Release]] %[[#Value]]
+
+  %10 = atomicrmw umin i32 addrspace(1)* @ui, i32 42 acq_rel
+; CHECK: %[[#]] = OpAtomicUMin %[[#Int]] %[[#Pointer]] %[[#Scope_Device]] %[[#MemSem_AcquireRelease]] %[[#Value]]
+
+  ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/exec_mode_float_control_khr.ll b/llvm/test/CodeGen/SPIRV/exec_mode_float_control_khr.ll
new file mode 100644
index 0000000000000..4da16528a8737
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/exec_mode_float_control_khr.ll
@@ -0,0 +1,72 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefixes=SPV
+
+define dso_local dllexport spir_kernel void @k_float_controls_0(i32 %ibuf, i32 %obuf) local_unnamed_addr {
+entry:
+  ret void
+}
+
+define dso_local dllexport spir_kernel void @k_float_controls_1(i32 %ibuf, i32 %obuf) local_unnamed_addr {
+entry:
+  ret void
+}
+
+define dso_local dllexport spir_kernel void @k_float_controls_2(i32 %ibuf, i32 %obuf) local_unnamed_addr {
+entry:
+  ret void
+}
+
+define dso_local dllexport spir_kernel void @k_float_controls_3(i32 %ibuf, i32 %obuf) local_unnamed_addr {
+entry:
+  ret void
+}
+
+define dso_local dllexport spir_kernel void @k_float_controls_4(i32 %ibuf, i32 %obuf) local_unnamed_addr {
+entry:
+  ret void
+}
+
+
+!spirv.ExecutionMode = !{!15, !16, !17, !18, !19, !20, !21, !22, !23, !24, !25, !26, !27, !28, !29}
+
+; SPV-NOT: OpExtension "SPV_KHR_float_controls"
+
+; SPV-DAG: OpEntryPoint {{.*}} %[[#KERNEL0:]] "k_float_controls_0"
+; SPV-DAG: OpEntryPoint {{.*}} %[[#KERNEL1:]] "k_float_controls_1"
+; SPV-DAG: OpEntryPoint {{.*}} %[[#KERNEL2:]] "k_float_controls_2"
+; SPV-DAG: OpEntryPoint {{.*}} %[[#KERNEL3:]] "k_float_controls_3"
+; SPV-DAG: OpEntryPoint {{.*}} %[[#KERNEL4:]] "k_float_controls_4"
+
+; SPV-DAG: OpExecutionMode %[[#KERNEL0]] DenormPreserve 64
+!15 = !{void (i32, i32)* @k_float_controls_0, i32 4459, i32 64}
+; SPV-DAG: OpExecutionMode %[[#KERNEL0]] DenormPreserve 32
+!16 = !{void (i32, i32)* @k_float_controls_0, i32 4459, i32 32}
+; SPV-DAG: OpExecutionMode %[[#KERNEL0]] DenormPreserve 16
+!17 = !{void (i32, i32)* @k_float_controls_0, i32 4459, i32 16}
+
+; SPV-DAG: OpExecutionMode %[[#KERNEL1]] DenormFlushToZero 64
+!18 = !{void (i32, i32)* @k_float_controls_1, i32 4460, i32 64}
+; SPV-DAG: OpExecutionMode %[[#KERNEL1]] DenormFlushToZero 32
+!19 = !{void (i32, i32)* @k_float_controls_1, i32 4460, i32 32}
+; SPV-DAG: OpExecutionMode %[[#KERNEL1]] DenormFlushToZero 16
+!20 = !{void (i32, i32)* @k_float_controls_1, i32 4460, i32 16}
+
+; SPV-DAG: OpExecutionMode %[[#KERNEL2]] SignedZeroInfNanPreserve 64
+!21 = !{void (i32, i32)* @k_float_controls_2, i32 4461, i32 64}
+; SPV-DAG: OpExecutionMode %[[#KERNEL2]] SignedZeroInfNanPreserve 32
+!22 = !{void (i32, i32)* @k_float_controls_2, i32 4461, i32 32}
+; SPV-DAG: OpExecutionMode %[[#KERNEL2]] SignedZeroInfNanPreserve 16
+!23 = !{void (i32, i32)* @k_float_controls_2, i32 4461, i32 16}
+
+; SPV-DAG: OpExecutionMode %[[#KERNEL3]] RoundingModeRTE 64
+!24 = !{void (i32, i32)* @k_float_controls_3, i32 4462, i32 64}
+; SPV-DAG: OpExecutionMode %[[#KERNEL3]] RoundingModeRTE 32
+!25 = !{void (i32, i32)* @k_float_controls_3, i32 4462, i32 32}
+; SPV-DAG: OpExecutionMode %[[#KERNEL3]] RoundingModeRTE 16
+!26 = !{void (i32, i32)* @k_float_controls_3, i32 4462, i32 16}
+
+; SPV-DAG: OpExecutionMode %[[#KERNEL4]] RoundingModeRTZ 64
+!27 = !{void (i32, i32)* @k_float_controls_4, i32 4463, i32 64}
+; SPV-DAG: OpExecutionMode %[[#KERNEL4]] RoundingModeRTZ 32
+!28 = !{void (i32, i32)* @k_float_controls_4, i32 4463, i32 32}
+; SPV-DAG: OpExecutionMode %[[#KERNEL4]] RoundingModeRTZ 16
+!29 = !{void (i32, i32)* @k_float_controls_4, i32 4463, i32 16}
diff --git a/llvm/test/CodeGen/SPIRV/instructions/call-trivial-function.ll b/llvm/test/CodeGen/SPIRV/instructions/call-trivial-function.ll
new file mode 100644
index 0000000000000..6924b7006f837
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/instructions/call-trivial-function.ll
@@ -0,0 +1,22 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
+
+; CHECK-DAG: OpName [[VALUE:%.+]] "value"
+; CHECK-DAG: OpName [[IDENTITY:%.+]] "identity"
+; CHECK-DAG: OpName [[FOO:%.+]] "foo"
+
+; CHECK:     [[INT:%.+]] = OpTypeInt 32
+; CHECK-DAG: [[CST:%.+]] = OpConstant [[INT]] 42
+
+define i32 @identity(i32 %value) {
+  ret i32 %value
+}
+
+define i32 @foo() {
+  %x = call i32 @identity(i32 42)
+  ret i32 %x
+}
+
+; CHECK: [[FOO]] = OpFunction [[INT]]
+; CHECK: [[X:%.+]] = OpFunctionCall [[INT]] [[IDENTITY]] [[CST]]
+; CHECK: OpReturnValue [[X]]
+; CHECK: OpFunctionEnd
diff --git a/llvm/test/CodeGen/SPIRV/instructions/scalar-bitwise-operations.ll b/llvm/test/CodeGen/SPIRV/instructions/scalar-bitwise-operations.ll
new file mode 100644
index 0000000000000..5424fb457bc7d
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/instructions/scalar-bitwise-operations.ll
@@ -0,0 +1,105 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
+
+; CHECK-DAG: OpName [[SCALAR_SHL:%.+]] "scalar_shl"
+; CHECK-DAG: OpName [[SCALAR_LSHR:%.+]] "scalar_lshr"
+; CHECK-DAG: OpName [[SCALAR_ASHR:%.+]] "scalar_ashr"
+; CHECK-DAG: OpName [[SCALAR_AND:%.+]] "scalar_and"
+; CHECK-DAG: OpName [[SCALAR_OR:%.+]] "scalar_or"
+; CHECK-DAG: OpName [[SCALAR_XOR:%.+]] "scalar_xor"
+
+; CHECK-NOT: DAG-FENCE
+
+; CHECK-DAG: [[SCALAR:%.+]] = OpTypeInt 32
+; CHECK-DAG: [[SCALAR_FN:%.+]] = OpTypeFunction [[SCALAR]] [[SCALAR]] [[SCALAR]]
+
+; CHECK-NOT: DAG-FENCE
+
+
+;; Test shl on scalar:
+define i32 @scalar_shl(i32 %a, i32 %b) {
+    %c = shl i32 %a, %b
+    ret i32 %c
+}
+
+; CHECK:      [[SCALAR_SHL]] = OpFunction [[SCALAR]] None [[SCALAR_FN]]
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[SCALAR]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[SCALAR]]
+; CHECK:      OpLabel
+; CHECK:      [[C:%.+]] = OpShiftLeftLogical [[SCALAR]] [[A]] [[B]]
+; CHECK:      OpReturnValue [[C]]
+; CHECK-NEXT: OpFunctionEnd
+
+
+;; Test lshr on scalar:
+define i32 @scalar_lshr(i32 %a, i32 %b) {
+    %c = lshr i32 %a, %b
+    ret i32 %c
+}
+
+; CHECK:      [[SCALAR_LSHR]] = OpFunction [[SCALAR]] None [[SCALAR_FN]]
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[SCALAR]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[SCALAR]]
+; CHECK:      OpLabel
+; CHECK:      [[C:%.+]] = OpShiftRightLogical [[SCALAR]] [[A]] [[B]]
+; CHECK:      OpReturnValue [[C]]
+; CHECK-NEXT: OpFunctionEnd
+
+
+;; Test ashr on scalar:
+define i32 @scalar_ashr(i32 %a, i32 %b) {
+    %c = ashr i32 %a, %b
+    ret i32 %c
+}
+
+; CHECK:      [[SCALAR_ASHR]] = OpFunction [[SCALAR]] None [[SCALAR_FN]]
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[SCALAR]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[SCALAR]]
+; CHECK:      OpLabel
+; CHECK:      [[C:%.+]] = OpShiftRightArithmetic [[SCALAR]] [[A]] [[B]]
+; CHECK:      OpReturnValue [[C]]
+; CHECK-NEXT: OpFunctionEnd
+
+
+;; Test and on scalar:
+define i32 @scalar_and(i32 %a, i32 %b) {
+    %c = and i32 %a, %b
+    ret i32 %c
+}
+
+; CHECK:      [[SCALAR_AND]] = OpFunction [[SCALAR]] None [[SCALAR_FN]]
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[SCALAR]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[SCALAR]]
+; CHECK:      OpLabel
+; CHECK:      [[C:%.+]] = OpBitwiseAnd [[SCALAR]] [[A]] [[B]]
+; CHECK:      OpReturnValue [[C]]
+; CHECK-NEXT: OpFunctionEnd
+
+
+;; Test or on scalar:
+define i32 @scalar_or(i32 %a, i32 %b) {
+    %c = or i32 %a, %b
+    ret i32 %c
+}
+
+; CHECK:      [[SCALAR_OR]] = OpFunction [[SCALAR]] None [[SCALAR_FN]]
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[SCALAR]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[SCALAR]]
+; CHECK:      OpLabel
+; CHECK:      [[C:%.+]] = OpBitwiseOr [[SCALAR]] [[A]] [[B]]
+; CHECK:      OpReturnValue [[C]]
+; CHECK-NEXT: OpFunctionEnd
+
+
+;; Test xor on scalar:
+define i32 @scalar_xor(i32 %a, i32 %b) {
+    %c = xor i32 %a, %b
+    ret i32 %c
+}
+
+; CHECK:      [[SCALAR_XOR]] = OpFunction [[SCALAR]] None [[SCALAR_FN]]
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[SCALAR]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[SCALAR]]
+; CHECK:      OpLabel
+; CHECK:      [[C:%.+]] = OpBitwiseXor [[SCALAR]] [[A]] [[B]]
+; CHECK:      OpReturnValue [[C]]
+; CHECK-NEXT: OpFunctionEnd
diff --git a/llvm/test/CodeGen/SPIRV/instructions/vector-bitwise-operations.ll b/llvm/test/CodeGen/SPIRV/instructions/vector-bitwise-operations.ll
new file mode 100644
index 0000000000000..664c42d805f63
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/instructions/vector-bitwise-operations.ll
@@ -0,0 +1,106 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
+
+; CHECK-DAG: OpName [[VECTOR_SHL:%.+]] "vector_shl"
+; CHECK-DAG: OpName [[VECTOR_LSHR:%.+]] "vector_lshr"
+; CHECK-DAG: OpName [[VECTOR_ASHR:%.+]] "vector_ashr"
+; CHECK-DAG: OpName [[VECTOR_AND:%.+]] "vector_and"
+; CHECK-DAG: OpName [[VECTOR_OR:%.+]] "vector_or"
+; CHECK-DAG: OpName [[VECTOR_XOR:%.+]] "vector_xor"
+
+; CHECK-NOT: DAG-FENCE
+
+; CHECK-DAG: [[I16:%.+]] = OpTypeInt 16
+; CHECK-DAG: [[VECTOR:%.+]] = OpTypeVector [[I16]]
+; CHECK-DAG: [[VECTOR_FN:%.+]] = OpTypeFunction [[VECTOR]] [[VECTOR]] [[VECTOR]]
+
+; CHECK-NOT: DAG-FENCE
+
+
+;; Test shl on vector:
+define <2 x i16> @vector_shl(<2 x i16> %a, <2 x i16> %b) {
+    %c = shl <2 x i16> %a, %b
+    ret <2 x i16> %c
+}
+
+; CHECK:      [[VECTOR_SHL]] = OpFunction [[VECTOR]] None [[VECTOR_FN]]
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[VECTOR]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[VECTOR]]
+; CHECK:      OpLabel
+; CHECK:      [[C:%.+]] = OpShiftLeftLogical [[VECTOR]] [[A]] [[B]]
+; CHECK:      OpReturnValue [[C]]
+; CHECK-NEXT: OpFunctionEnd
+
+
+;; Test lshr on vector:
+define <2 x i16> @vector_lshr(<2 x i16> %a, <2 x i16> %b) {
+    %c = lshr <2 x i16> %a, %b
+    ret <2 x i16> %c
+}
+
+; CHECK:      [[VECTOR_LSHR]] = OpFunction [[VECTOR]] None [[VECTOR_FN]]
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[VECTOR]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[VECTOR]]
+; CHECK:      OpLabel
+; CHECK:      [[C:%.+]] = OpShiftRightLogical [[VECTOR]] [[A]] [[B]]
+; CHECK:      OpReturnValue [[C]]
+; CHECK-NEXT: OpFunctionEnd
+
+
+;; Test ashr on vector:
+define <2 x i16> @vector_ashr(<2 x i16> %a, <2 x i16> %b) {
+    %c = ashr <2 x i16> %a, %b
+    ret <2 x i16> %c
+}
+
+; CHECK:      [[VECTOR_ASHR]] = OpFunction [[VECTOR]] None [[VECTOR_FN]]
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[VECTOR]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[VECTOR]]
+; CHECK:      OpLabel
+; CHECK:      [[C:%.+]] = OpShiftRightArithmetic [[VECTOR]] [[A]] [[B]]
+; CHECK:      OpReturnValue [[C]]
+; CHECK-NEXT: OpFunctionEnd
+
+
+;; Test and on vector:
+define <2 x i16> @vector_and(<2 x i16> %a, <2 x i16> %b) {
+    %c = and <2 x i16> %a, %b
+    ret <2 x i16> %c
+}
+
+; CHECK:      [[VECTOR_AND]] = OpFunction [[VECTOR]] None [[VECTOR_FN]]
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[VECTOR]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[VECTOR]]
+; CHECK:      OpLabel
+; CHECK:      [[C:%.+]] = OpBitwiseAnd [[VECTOR]] [[A]] [[B]]
+; CHECK:      OpReturnValue [[C]]
+; CHECK-NEXT: OpFunctionEnd
+
+
+;; Test or on vector:
+define <2 x i16> @vector_or(<2 x i16> %a, <2 x i16> %b) {
+    %c = or <2 x i16> %a, %b
+    ret <2 x i16> %c
+}
+
+; CHECK:      [[VECTOR_OR]] = OpFunction [[VECTOR]] None [[VECTOR_FN]]
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[VECTOR]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[VECTOR]]
+; CHECK:      OpLabel
+; CHECK:      [[C:%.+]] = OpBitwiseOr [[VECTOR]] [[A]] [[B]]
+; CHECK:      OpReturnValue [[C]]
+; CHECK-NEXT: OpFunctionEnd
+
+
+;; Test xor on vector:
+define <2 x i16> @vector_xor(<2 x i16> %a, <2 x i16> %b) {
+    %c = xor <2 x i16> %a, %b
+    ret <2 x i16> %c
+}
+
+; CHECK:      [[VECTOR_XOR]] = OpFunction [[VECTOR]] None [[VECTOR_FN]]
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[VECTOR]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[VECTOR]]
+; CHECK:      OpLabel
+; CHECK:      [[C:%.+]] = OpBitwiseXor [[VECTOR]] [[A]] [[B]]
+; CHECK:      OpReturnValue [[C]]
+; CHECK-NEXT: OpFunctionEnd
diff --git a/llvm/test/CodeGen/SPIRV/instructions/vector-shuffle.ll b/llvm/test/CodeGen/SPIRV/instructions/vector-shuffle.ll
new file mode 100644
index 0000000000000..2c5d528df0fc4
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/instructions/vector-shuffle.ll
@@ -0,0 +1,68 @@
+; RUN: llc -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
+
+; CHECK-DAG: OpName [[SHFv4:%.+]] "shuffle_v4"
+; CHECK-DAG: OpName [[INSv4:%.+]] "insert_v4"
+; CHECK-DAG: OpName [[EXTv4:%.+]] "extract_v4"
+; CHECK-DAG: OpName [[INSv4C:%.+]] "insert_v4C"
+; CHECK-DAG: OpName [[EXTv4C:%.+]] "extract_v4C"
+
+
+; CHECK:      [[SHFv4]] = OpFunction
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter
+; CHECK:      OpLabel
+; CHECK:      [[R:%.+]] = OpVectorShuffle {{%.+}} [[A]] [[B]] 0 4 3 6
+; CHECK:      OpReturnValue [[R]]
+; CHECK-NEXT: OpFunctionEnd
+define <4 x float> @shuffle_v4(<8 x float> %A, <8 x float> %B) {
+  %r = shufflevector <8 x float> %A, <8 x float> %B, <4 x i32> <i32 0, i32 4, i32 3, i32 6>
+  ret <4 x float> %r
+}
+
+; CHECK:      [[INSv4]] = OpFunction
+; CHECK-NEXT: [[V:%.+]] = OpFunctionParameter
+; CHECK-NEXT: [[E:%.+]] = OpFunctionParameter
+; CHECK-NEXT: [[C:%.+]] = OpFunctionParameter
+; CHECK:      OpLabel
+; CHECK:      [[R:%.+]] = OpVectorInsertDynamic {{%.+}} [[V]] [[E]] [[C]]
+; CHECK:      OpReturnValue [[R]]
+; CHECK-NEXT: OpFunctionEnd
+define <4 x float> @insert_v4(<4 x float> %V, float %E, i32 %C) {
+  %r = insertelement <4 x float> %V, float %E, i32 %C
+  ret <4 x float> %r
+}
+
+; CHECK:      [[EXTv4]] = OpFunction
+; CHECK-NEXT: [[V:%.+]] = OpFunctionParameter
+; CHECK-NEXT: [[C:%.+]] = OpFunctionParameter
+; CHECK:      OpLabel
+; CHECK:      [[R:%.+]] = OpVectorExtractDynamic {{%.+}} [[V]] [[C]]
+; CHECK:      OpReturnValue [[R]]
+; CHECK-NEXT: OpFunctionEnd
+define float @extract_v4(<4 x float> %V, i32 %C) {
+  %r = extractelement <4 x float> %V, i32 %C
+  ret float %r
+}
+
+; CHECK:      [[INSv4C]] = OpFunction
+; CHECK-NEXT: [[V:%.+]] = OpFunctionParameter
+; CHECK-NEXT: [[E:%.+]] = OpFunctionParameter
+; CHECK:      OpLabel
+; CHECK:      [[R:%.+]] = OpCompositeInsert {{%.+}} [[E]] [[V]] 3
+; CHECK:      OpReturnValue [[R]]
+; CHECK-NEXT: OpFunctionEnd
+define <4 x float> @insert_v4C(<4 x float> %V, float %E) {
+  %r = insertelement <4 x float> %V, float %E, i32 3
+  ret <4 x float> %r
+}
+
+; CHECK:      [[EXTv4C]] = OpFunction
+; CHECK-NEXT: [[V:%.+]] = OpFunctionParameter
+; CHECK:      OpLabel
+; CHECK:      [[R:%.+]] = OpCompositeExtract {{%.+}} [[V]] 2
+; CHECK:      OpReturnValue [[R]]
+; CHECK-NEXT: OpFunctionEnd
+define float @extract_v4C(<4 x float> %V) {
+  %r = extractelement <4 x float> %V, i32 2
+  ret float %r
+}
diff --git a/llvm/test/CodeGen/SPIRV/linked-list.ll b/llvm/test/CodeGen/SPIRV/linked-list.ll
new file mode 100644
index 0000000000000..126b030d64bc2
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/linked-list.ll
@@ -0,0 +1,10 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
+
+%struct.Node = type { %struct.Node.0 addrspace(1)* }
+; CHECK: %[[#]] = OpTypeOpaque "struct.Node.0"
+%struct.Node.0 = type opaque
+
+define spir_kernel void @create_linked_lists(%struct.Node addrspace(1)* nocapture %pNodes, i32 addrspace(1)* nocapture %allocation_index, i32 %list_length) {
+entry:
+  ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/assume.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/assume.ll
new file mode 100644
index 0000000000000..48c96fae8b03a
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/assume.ll
@@ -0,0 +1,56 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV-NOT: OpCapability ExpectAssumeKHR
+; CHECK-SPIRV-NOT: OpExtension "SPV_KHR_expect_assume"
+; CHECK-SPIRV:     OpName %[[#COMPARE:]] "cmp"
+; CHECK-SPIRV:     %[[#COMPARE]] = OpINotEqual %[[#]] %[[#]] %[[#]]
+; CHECK-SPIRV-NOT: OpAssumeTrueKHR %[[#COMPARE]]
+
+%class.anon = type { i8 }
+
+define spir_func void @_Z3fooi(i32 %x) {
+entry:
+  %x.addr = alloca i32, align 4
+  store i32 %x, i32* %x.addr, align 4
+  %0 = load i32, i32* %x.addr, align 4
+  %cmp = icmp ne i32 %0, 0
+  call void @llvm.assume(i1 %cmp)
+  ret void
+}
+
+declare void @llvm.assume(i1)
+
+define i32 @main() {
+entry:
+  %retval = alloca i32, align 4
+  %agg.tmp = alloca %class.anon, align 1
+  store i32 0, i32* %retval, align 4
+  call spir_func void @"_Z18kernel_single_taskIZ4mainE11fake_kernelZ4mainE3$_0EvT0_"(%class.anon* byval(%class.anon) align 1 %agg.tmp)
+  ret i32 0
+}
+
+define internal spir_func void @"_Z18kernel_single_taskIZ4mainE11fake_kernelZ4mainE3$_0EvT0_"(%class.anon* byval(%class.anon) align 1 %kernelFunc) {
+entry:
+  call spir_func void @"_ZZ4mainENK3$_0clEv"(%class.anon* %kernelFunc)
+  ret void
+}
+
+define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%class.anon* %this) align 2 {
+entry:
+  %this.addr = alloca %class.anon*, align 8
+  %a = alloca i32, align 4
+  store %class.anon* %this, %class.anon** %this.addr, align 8
+  %this1 = load %class.anon*, %class.anon** %this.addr, align 8
+  %0 = bitcast i32* %a to i8*
+  call void @llvm.lifetime.start.p0i8(i64 4, i8* %0)
+  store i32 1, i32* %a, align 4
+  %1 = load i32, i32* %a, align 4
+  call spir_func void @_Z3fooi(i32 %1)
+  %2 = bitcast i32* %a to i8*
+  call void @llvm.lifetime.end.p0i8(i64 4, i8* %2)
+  ret void
+}
+
+declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture)
+
+declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture)
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/expect.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/expect.ll
new file mode 100644
index 0000000000000..ec40c263be411
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/expect.ll
@@ -0,0 +1,109 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV-NOT: OpCapability ExpectAssumeKHR
+; CHECK-SPIRV-NOT: OpExtension "SPV_KHR_expect_assume"
+; CHECK-SPIRV:     OpFunction
+; CHECK-SPIRV-NOT: %[[#]] = OpExpectKHR %[[#]] %[[#]] %[[#]]
+; CHECK-SPIRV:     %[[#RES1:]] = OpSConvert %[[#]] %[[#]]
+; CHECK-SPIRV:     %[[#]] = OpINotEqual %[[#]] %[[#RES1]] %[[#]]
+
+; CHECK-SPIRV:     OpFunction
+; CHECK-SPIRV:     %[[#RES2:]] = OpSConvert %[[#]] %[[#]]
+; CHECK-SPIRV-NOT: %[[#]] = OpExpectKHR %[[#]] %[[#]] %[[#]]
+; CHECK-SPIRV:     %[[#]] = OpINotEqual %[[#]] %[[#RES2]] %[[#]]
+
+%"class._ZTSZ4mainE3$_0.anon" = type { i8 }
+
+define spir_kernel void @_ZTSZ4mainE15kernel_function() {
+entry:
+  %0 = alloca %"class._ZTSZ4mainE3$_0.anon", align 1
+  %1 = bitcast %"class._ZTSZ4mainE3$_0.anon"* %0 to i8*
+  call void @llvm.lifetime.start.p0i8(i64 1, i8* %1)
+  %2 = addrspacecast %"class._ZTSZ4mainE3$_0.anon"* %0 to %"class._ZTSZ4mainE3$_0.anon" addrspace(4)*
+  call spir_func void @"_ZZ4mainENK3$_0clEv"(%"class._ZTSZ4mainE3$_0.anon" addrspace(4)* %2)
+  %3 = bitcast %"class._ZTSZ4mainE3$_0.anon"* %0 to i8*
+  call void @llvm.lifetime.end.p0i8(i64 1, i8* %3)
+  ret void
+}
+
+declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture)
+
+define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%"class._ZTSZ4mainE3$_0.anon" addrspace(4)* %this) align 2 {
+entry:
+  %this.addr = alloca %"class._ZTSZ4mainE3$_0.anon" addrspace(4)*, align 8
+  %a = alloca i32, align 4
+  %b = alloca i32, align 4
+  store %"class._ZTSZ4mainE3$_0.anon" addrspace(4)* %this, %"class._ZTSZ4mainE3$_0.anon" addrspace(4)** %this.addr, align 8
+  %this1 = load %"class._ZTSZ4mainE3$_0.anon" addrspace(4)*, %"class._ZTSZ4mainE3$_0.anon" addrspace(4)** %this.addr, align 8
+  %0 = bitcast i32* %a to i8*
+  call void @llvm.lifetime.start.p0i8(i64 4, i8* %0)
+  %call = call spir_func i32 @_Z12expect_consti(i32 1)
+  store i32 %call, i32* %a, align 4
+  %1 = bitcast i32* %b to i8*
+  call void @llvm.lifetime.start.p0i8(i64 4, i8* %1)
+  %call2 = call spir_func i32 @_Z10expect_funi(i32 2)
+  store i32 %call2, i32* %b, align 4
+  %2 = bitcast i32* %b to i8*
+  call void @llvm.lifetime.end.p0i8(i64 4, i8* %2)
+  %3 = bitcast i32* %a to i8*
+  call void @llvm.lifetime.end.p0i8(i64 4, i8* %3)
+  ret void
+}
+
+declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture)
+
+define spir_func i32 @_Z12expect_consti(i32 %x) {
+entry:
+  %retval = alloca i32, align 4
+  %x.addr = alloca i32, align 4
+  store i32 %x, i32* %x.addr, align 4
+  %0 = load i32, i32* %x.addr, align 4
+  %conv = sext i32 %0 to i64
+  %expval = call i64 @llvm.expect.i64(i64 %conv, i64 1)
+  %tobool = icmp ne i64 %expval, 0
+  br i1 %tobool, label %if.then, label %if.end
+
+if.then:                                          ; preds = %entry
+  store i32 0, i32* %retval, align 4
+  br label %return
+
+if.end:                                           ; preds = %entry
+  %1 = load i32, i32* %x.addr, align 4
+  store i32 %1, i32* %retval, align 4
+  br label %return
+
+return:                                           ; preds = %if.end, %if.then
+  %2 = load i32, i32* %retval, align 4
+  ret i32 %2
+}
+
+define spir_func i32 @_Z10expect_funi(i32 %x) {
+entry:
+  %retval = alloca i32, align 4
+  %x.addr = alloca i32, align 4
+  store i32 %x, i32* %x.addr, align 4
+  %0 = load i32, i32* %x.addr, align 4
+  %conv = sext i32 %0 to i64
+  %call = call spir_func i32 @_Z3foov()
+  %conv1 = sext i32 %call to i64
+  %expval = call i64 @llvm.expect.i64(i64 %conv, i64 %conv1)
+  %tobool = icmp ne i64 %expval, 0
+  br i1 %tobool, label %if.then, label %if.end
+
+if.then:                                          ; preds = %entry
+  store i32 0, i32* %retval, align 4
+  br label %return
+
+if.end:                                           ; preds = %entry
+  %1 = load i32, i32* %x.addr, align 4
+  store i32 %1, i32* %retval, align 4
+  br label %return
+
+return:                                           ; preds = %if.end, %if.then
+  %2 = load i32, i32* %retval, align 4
+  ret i32 %2
+}
+
+declare i64 @llvm.expect.i64(i64, i64)
+
+declare spir_func i32 @_Z3foov()
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/fshr.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/fshr.ll
new file mode 100644
index 0000000000000..4cf5ca53a411b
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/fshr.ll
@@ -0,0 +1,85 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV:     OpName %[[#NAME_FSHR_FUNC_32:]] "spirv.llvm_fshr_i32"
+; CHECK-SPIRV:     OpName %[[#NAME_FSHR_FUNC_16:]] "spirv.llvm_fshr_i16"
+; CHECK-SPIRV:     OpName %[[#NAME_FSHR_FUNC_VEC_INT_16:]] "spirv.llvm_fshr_v2i16"
+; CHECK-SPIRV:     %[[#TYPE_INT_32:]] = OpTypeInt 32 0
+; CHECK-SPIRV:     %[[#TYPE_ORIG_FUNC_32:]] = OpTypeFunction %[[#TYPE_INT_32]] %[[#TYPE_INT_32]] %[[#TYPE_INT_32]]
+; CHECK-SPIRV:     %[[#TYPE_INT_16:]] = OpTypeInt 16 0
+; CHECK-SPIRV:     %[[#TYPE_ORIG_FUNC_16:]] = OpTypeFunction %[[#TYPE_INT_16]] %[[#TYPE_INT_16]] %[[#TYPE_INT_16]]
+; CHECK-SPIRV:     %[[#TYPE_VEC_INT_16:]] = OpTypeVector %[[#TYPE_INT_16]] 2
+; CHECK-SPIRV:     %[[#TYPE_ORIG_FUNC_VEC_INT_16:]] = OpTypeFunction %[[#TYPE_VEC_INT_16]] %[[#TYPE_VEC_INT_16]] %[[#TYPE_VEC_INT_16]]
+; CHECK-SPIRV:     %[[#TYPE_FSHR_FUNC_32:]] = OpTypeFunction %[[#TYPE_INT_32]] %[[#TYPE_INT_32]] %[[#TYPE_INT_32]] %[[#TYPE_INT_32]]
+; CHECK-SPIRV:     %[[#TYPE_FSHR_FUNC_16:]] = OpTypeFunction %[[#TYPE_INT_16]] %[[#TYPE_INT_16]] %[[#TYPE_INT_16]] %[[#TYPE_INT_16]]
+; CHECK-SPIRV:     %[[#TYPE_FSHR_FUNC_VEC_INT_16:]] = OpTypeFunction %[[#TYPE_VEC_INT_16]] %[[#TYPE_VEC_INT_16]] %[[#TYPE_VEC_INT_16]] %[[#TYPE_VEC_INT_16]]
+; CHECK-SPIRV-DAG: %[[#CONST_ROTATE_32:]] = OpConstant %[[#TYPE_INT_32]] 8
+; CHECK-SPIRV-DAG: %[[#CONST_ROTATE_16:]] = OpConstant %[[#TYPE_INT_16]] 8
+; CHECK-SPIRV:     %[[#CONST_ROTATE_VEC_INT_16:]] = OpConstantComposite %[[#TYPE_VEC_INT_16]] %[[#CONST_ROTATE_16]] %[[#CONST_ROTATE_16]]
+; CHECK-SPIRV-DAG: %[[#CONST_TYPE_SIZE_32:]] = OpConstant %[[#TYPE_INT_32]] 32
+
+; CHECK-SPIRV: %[[#]] = OpFunction %[[#TYPE_INT_32]] {{.*}} %[[#TYPE_ORIG_FUNC_32]]
+; CHECK-SPIRV: %[[#X:]] = OpFunctionParameter %[[#TYPE_INT_32]]
+; CHECK-SPIRV: %[[#Y:]] = OpFunctionParameter %[[#TYPE_INT_32]]
+define spir_func i32 @Test_i32(i32 %x, i32 %y) local_unnamed_addr {
+entry:
+  ; CHECK-SPIRV: %[[#CALL_32_X_Y:]] = OpFunctionCall %[[#TYPE_INT_32]] %[[#NAME_FSHR_FUNC_32]] %[[#X]] %[[#Y]] %[[#CONST_ROTATE_32]]
+  %0 = call i32 @llvm.fshr.i32(i32 %x, i32 %y, i32 8)
+  ; CHECK-SPIRV: %[[#CALL_32_Y_X:]] = OpFunctionCall %[[#TYPE_INT_32]] %[[#NAME_FSHR_FUNC_32]] %[[#Y]] %[[#X]] %[[#CONST_ROTATE_32]]
+  %1 = call i32 @llvm.fshr.i32(i32 %y, i32 %x, i32 8)
+  ; CHECK-SPIRV: %[[#ADD_32:]] = OpIAdd %[[#TYPE_INT_32]] %[[#CALL_32_X_Y]] %[[#CALL_32_Y_X]]
+  %sum = add i32 %0, %1
+  ; CHECK-SPIRV: OpReturnValue %[[#ADD_32]]
+  ret i32 %sum
+}
+
+; CHECK-SPIRV: %[[#]] = OpFunction %[[#TYPE_INT_16]] {{.*}} %[[#TYPE_ORIG_FUNC_16]]
+; CHECK-SPIRV: %[[#X:]] = OpFunctionParameter %[[#TYPE_INT_16]]
+; CHECK-SPIRV: %[[#Y:]] = OpFunctionParameter %[[#TYPE_INT_16]]
+define spir_func i16 @Test_i16(i16 %x, i16 %y) local_unnamed_addr {
+entry:
+  ; CHECK-SPIRV: %[[#CALL_16:]] = OpFunctionCall %[[#TYPE_INT_16]] %[[#NAME_FSHR_FUNC_16]] %[[#X]] %[[#Y]] %[[#CONST_ROTATE_16]]
+  %0 = call i16 @llvm.fshr.i16(i16 %x, i16 %y, i16 8)
+  ; CHECK-SPIRV: OpReturnValue %[[#CALL_16]]
+  ret i16 %0
+}
+
+; CHECK-SPIRV: %[[#]] = OpFunction %[[#TYPE_VEC_INT_16]] {{.*}} %[[#TYPE_ORIG_FUNC_VEC_INT_16]]
+; CHECK-SPIRV: %[[#X:]] = OpFunctionParameter %[[#TYPE_VEC_INT_16]]
+; CHECK-SPIRV: %[[#Y:]] = OpFunctionParameter %[[#TYPE_VEC_INT_16]]
+define spir_func <2 x i16> @Test_v2i16(<2 x i16> %x, <2 x i16> %y) local_unnamed_addr {
+entry:
+  ; CHECK-SPIRV: %[[#CALL_VEC_INT_16:]] = OpFunctionCall %[[#TYPE_VEC_INT_16]] %[[#NAME_FSHR_FUNC_VEC_INT_16]] %[[#X]] %[[#Y]] %[[#CONST_ROTATE_VEC_INT_16]]
+  %0 = call <2 x i16> @llvm.fshr.v2i16(<2 x i16> %x, <2 x i16> %y, <2 x i16> <i16 8, i16 8>)
+  ; CHECK-SPIRV: OpReturnValue %[[#CALL_VEC_INT_16]]
+  ret <2 x i16> %0
+}
+
+; CHECK-SPIRV: %[[#NAME_FSHR_FUNC_32]] = OpFunction %[[#TYPE_INT_32]] {{.*}} %[[#TYPE_FSHR_FUNC_32]]
+; CHECK-SPIRV: %[[#X_ARG:]] = OpFunctionParameter %[[#TYPE_INT_32]]
+; CHECK-SPIRV: %[[#Y_ARG:]] = OpFunctionParameter %[[#TYPE_INT_32]]
+; CHECK-SPIRV: %[[#ROT:]] = OpFunctionParameter %[[#TYPE_INT_32]]
+
+; CHECK-SPIRV: %[[#ROTATE_MOD_SIZE:]] = OpUMod %[[#TYPE_INT_32]] %[[#ROT]] %[[#CONST_TYPE_SIZE_32]]
+; CHECK-SPIRV: %[[#Y_SHIFT_RIGHT:]] = OpShiftRightLogical %[[#TYPE_INT_32]] %[[#Y_ARG]] %[[#ROTATE_MOD_SIZE]]
+; CHECK-SPIRV: %[[#NEG_ROTATE:]] = OpISub %[[#TYPE_INT_32]] %[[#CONST_TYPE_SIZE_32]] %[[#ROTATE_MOD_SIZE]]
+; CHECK-SPIRV: %[[#X_SHIFT_LEFT:]] = OpShiftLeftLogical %[[#TYPE_INT_32]] %[[#X_ARG]] %[[#NEG_ROTATE]]
+; CHECK-SPIRV: %[[#FSHR_RESULT:]] = OpBitwiseOr %[[#TYPE_INT_32]] %[[#Y_SHIFT_RIGHT]] %[[#X_SHIFT_LEFT]]
+; CHECK-SPIRV: OpReturnValue %[[#FSHR_RESULT]]
+
+;; Just check that the function for i16 was generated as such - we've checked the logic for another type.
+; CHECK-SPIRV: %[[#NAME_FSHR_FUNC_16]] = OpFunction %[[#TYPE_INT_16]] {{.*}} %[[#TYPE_FSHR_FUNC_16]]
+; CHECK-SPIRV: %[[#X_ARG:]] = OpFunctionParameter %[[#TYPE_INT_16]]
+; CHECK-SPIRV: %[[#Y_ARG:]] = OpFunctionParameter %[[#TYPE_INT_16]]
+; CHECK-SPIRV: %[[#ROT:]] = OpFunctionParameter %[[#TYPE_INT_16]]
+
+;; Just check that the function for v2i16 was generated as such - we've checked the logic for another type.
+; CHECK-SPIRV: %[[#NAME_FSHR_FUNC_VEC_INT_16]] = OpFunction %[[#TYPE_VEC_INT_16]] {{.*}} %[[#TYPE_FSHR_FUNC_VEC_INT_16]]
+; CHECK-SPIRV: %[[#X_ARG:]] = OpFunctionParameter %[[#TYPE_VEC_INT_16]]
+; CHECK-SPIRV: %[[#Y_ARG:]] = OpFunctionParameter %[[#TYPE_VEC_INT_16]]
+; CHECK-SPIRV: %[[#ROT:]] = OpFunctionParameter %[[#TYPE_VEC_INT_16]]
+
+declare i32 @llvm.fshr.i32(i32, i32, i32)
+
+declare i16 @llvm.fshr.i16(i16, i16, i16)
+
+declare <2 x i16> @llvm.fshr.v2i16(<2 x i16>, <2 x i16>, <2 x i16>)
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/invariant.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/invariant.ll
new file mode 100644
index 0000000000000..5b700b7fcfee3
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/invariant.ll
@@ -0,0 +1,19 @@
+;; Make sure the backend doesn't crash if the input LLVM IR contains llvm.invariant.* intrinsics
+; RUN: llc -O0 -mtriple=spirv64-unknown-linux %s -o - | FileCheck %s
+
+; CHECK-NOT: OpFunctionParameter
+; CHECK-NOT: OpFunctionCall
+
+@WGSharedVar = internal addrspace(3) constant i64 0, align 8
+
+declare {}* @llvm.invariant.start.p3i8(i64 immarg, i8 addrspace(3)* nocapture)
+
+declare void @llvm.invariant.end.p3i8({}*, i64 immarg, i8 addrspace(3)* nocapture)
+
+define linkonce_odr dso_local spir_func void @func() {
+  store i64 2, i64 addrspace(3)* @WGSharedVar
+  %1 = bitcast i64 addrspace(3)* @WGSharedVar to i8 addrspace(3)*
+  %2 = call {}* @llvm.invariant.start.p3i8(i64 8, i8 addrspace(3)* %1)
+  call void @llvm.invariant.end.p3i8({}* %2, i64 8, i8 addrspace(3)* %1)
+  ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/umul.with.overflow.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/umul.with.overflow.ll
new file mode 100644
index 0000000000000..406a23fa7d3df
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/umul.with.overflow.ll
@@ -0,0 +1,54 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: OpName %[[#NAME_UMUL_FUNC_8:]] "spirv.llvm_umul_with_overflow_i8"
+; CHECK-SPIRV: OpName %[[#NAME_UMUL_FUNC_32:]] "spirv.llvm_umul_with_overflow_i32"
+; CHECK-SPIRV: OpName %[[#NAME_UMUL_FUNC_VEC_I64:]] "spirv.llvm_umul_with_overflow_v2i64"
+
+define dso_local spir_func void @_Z4foo8hhPh(i8 zeroext %a, i8 zeroext %b, i8* nocapture %c) local_unnamed_addr {
+entry:
+  ; CHECK-SPIRV: %[[#]] = OpFunctionCall %[[#]] %[[#NAME_UMUL_FUNC_8]]
+  %umul = tail call { i8, i1 } @llvm.umul.with.overflow.i8(i8 %a, i8 %b)
+  %cmp = extractvalue { i8, i1 } %umul, 1
+  %umul.value = extractvalue { i8, i1 } %umul, 0
+  %storemerge = select i1 %cmp, i8 0, i8 %umul.value
+  store i8 %storemerge, i8* %c, align 1
+  ret void
+}
+
+define dso_local spir_func void @_Z5foo32jjPj(i32 %a, i32 %b, i32* nocapture %c) local_unnamed_addr {
+entry:
+  ; CHECK-SPIRV: %[[#]] = OpFunctionCall %[[#]] %[[#NAME_UMUL_FUNC_32]]
+  %umul = tail call { i32, i1 } @llvm.umul.with.overflow.i32(i32 %b, i32 %a)
+  %umul.val = extractvalue { i32, i1 } %umul, 0
+  %umul.ov = extractvalue { i32, i1 } %umul, 1
+  %spec.select = select i1 %umul.ov, i32 0, i32 %umul.val
+  store i32 %spec.select, i32* %c, align 4
+  ret void
+}
+
+define dso_local spir_func void @umulo_v2i64(<2 x i64> %a, <2 x i64> %b, <2 x i64>* %p) nounwind {
+  ; CHECK-SPIRV: %[[#]] = OpFunctionCall %[[#]] %[[#NAME_UMUL_FUNC_VEC_I64]]
+  %umul = call {<2 x i64>, <2 x i1>} @llvm.umul.with.overflow.v2i64(<2 x i64> %a, <2 x i64> %b)
+  %umul.val = extractvalue {<2 x i64>, <2 x i1>} %umul, 0
+  %umul.ov = extractvalue {<2 x i64>, <2 x i1>} %umul, 1
+  %zero = alloca <2 x i64>, align 16
+  %spec.select = select <2 x i1> %umul.ov, <2 x i64> <i64 0, i64 0>, <2 x i64> %umul.val
+  store <2 x i64> %spec.select, <2 x i64>* %p
+  ret void
+}
+
+; CHECK-SPIRV: %[[#NAME_UMUL_FUNC_8]] = OpFunction %[[#]]
+; CHECK-SPIRV: %[[#VAR_A:]] = OpFunctionParameter %[[#]]
+; CHECK-SPIRV: %[[#VAR_B:]] = OpFunctionParameter %[[#]]
+; CHECK-SPIRV: %[[#MUL_RES:]] = OpIMul %[[#]] %[[#VAR_A]] %[[#VAR_B]]
+; CHECK-SPIRV: %[[#DIV_RES:]] = OpUDiv %[[#]] %[[#MUL_RES]] %[[#VAR_A]]
+; CHECK-SPIRV: %[[#CMP_RES:]] = OpINotEqual %[[#]] %[[#VAR_A]] %[[#DIV_RES]]
+; CHECK-SPIRV: %[[#INSERT_RES:]] = OpCompositeInsert %[[#]] %[[#MUL_RES]]
+; CHECK-SPIRV: %[[#INSERT_RES_1:]] = OpCompositeInsert %[[#]] %[[#CMP_RES]] %[[#INSERT_RES]]
+; CHECK-SPIRV: OpReturnValue %[[#INSERT_RES_1]]
+
+declare { i8, i1 } @llvm.umul.with.overflow.i8(i8, i8)
+
+declare { i32, i1 } @llvm.umul.with.overflow.i32(i32, i32)
+
+declare {<2 x i64>, <2 x i1>} @llvm.umul.with.overflow.v2i64(<2 x i64>, <2 x i64>)
diff --git a/llvm/test/CodeGen/SPIRV/memory_model_md.ll b/llvm/test/CodeGen/SPIRV/memory_model_md.ll
new file mode 100644
index 0000000000000..e52343cbbb7e4
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/memory_model_md.ll
@@ -0,0 +1,11 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPV
+
+; SPV: OpMemoryModel Physical32 Simple
+define dso_local dllexport void @k_no_fc(i32 %ibuf, i32 %obuf) local_unnamed_addr {
+entry:
+  ret void
+}
+
+!spirv.MemoryModel = !{!0}
+
+!0 = !{i32 1, i32 0}
diff --git a/llvm/test/CodeGen/SPIRV/multi_md.ll b/llvm/test/CodeGen/SPIRV/multi_md.ll
new file mode 100644
index 0000000000000..6d8af7defd752
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/multi_md.ll
@@ -0,0 +1,50 @@
+;; Check duplicate operands in opencl.ocl.version metadata is accepted without
+;; assertion.
+
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+
+%struct.my_struct_t = type { i8, i32 }
+
+@var = addrspace(1) global %struct.my_struct_t { i8 97, i32 42 }, align 4
+
+define spir_kernel void @__OpenCL_writer_kernel(i8 zeroext %c, i32 %i) {
+entry:
+  %c.addr = alloca i8, align 1
+  %i.addr = alloca i32, align 4
+  store i8 %c, i8* %c.addr, align 1
+  store i32 %i, i32* %i.addr, align 4
+  %0 = load i8, i8* %c.addr, align 1
+  store i8 %0, i8 addrspace(1)* getelementptr inbounds (%struct.my_struct_t, %struct.my_struct_t addrspace(1)* @var, i32 0, i32 0), align 1
+  %1 = load i32, i32* %i.addr, align 4
+  store i32 %1, i32 addrspace(1)* getelementptr inbounds (%struct.my_struct_t, %struct.my_struct_t addrspace(1)* @var, i32 0, i32 1), align 4
+  ret void
+}
+
+define spir_kernel void @__OpenCL_reader_kernel(i8 addrspace(1)* %C, i32 addrspace(1)* %I) {
+entry:
+  %C.addr = alloca i8 addrspace(1)*, align 8
+  %I.addr = alloca i32 addrspace(1)*, align 8
+  store i8 addrspace(1)* %C, i8 addrspace(1)** %C.addr, align 8
+  store i32 addrspace(1)* %I, i32 addrspace(1)** %I.addr, align 8
+  %0 = load i8, i8 addrspace(1)* getelementptr inbounds (%struct.my_struct_t, %struct.my_struct_t addrspace(1)* @var, i32 0, i32 0), align 1
+  %1 = load i8 addrspace(1)*, i8 addrspace(1)** %C.addr, align 8
+  store i8 %0, i8 addrspace(1)* %1, align 1
+  %2 = load i32, i32 addrspace(1)* getelementptr inbounds (%struct.my_struct_t, %struct.my_struct_t addrspace(1)* @var, i32 0, i32 1), align 4
+  %3 = load i32 addrspace(1)*, i32 addrspace(1)** %I.addr, align 8
+  store i32 %2, i32 addrspace(1)* %3, align 4
+  ret void
+}
+
+;; "cl_images" should be encoded as BasicImage capability,
+;; but images are not used in this test case, so this capability is not required.
+; CHECK-NOT: OpExtension "cl_images"
+; CHECK-DAG: OpSourceExtension "cl_khr_int64_base_atomics"
+; CHECK-DAG: OpSourceExtension "cl_khr_int64_extended_atomics"
+; CHECK:     OpSource OpenCL_C 200000
+
+!opencl.ocl.version = !{!13, !13}
+!opencl.used.extensions = !{!24, !25}
+
+!13 = !{i32 2, i32 0}
+!24 = !{!"cl_khr_int64_base_atomics"}
+!25 = !{!"cl_khr_int64_base_atomics", !"cl_khr_int64_extended_atomics"}
diff --git a/llvm/test/CodeGen/SPIRV/opencl/basic/get_global_offset.ll b/llvm/test/CodeGen/SPIRV/opencl/basic/get_global_offset.ll
new file mode 100644
index 0000000000000..2d07dae0a63f1
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/opencl/basic/get_global_offset.ll
@@ -0,0 +1,57 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+
+; CHECK: OpEntryPoint Kernel %[[#test_func:]] "test"
+; CHECK: OpName %[[#outOffsets:]] "outOffsets"
+; CHECK: OpName %[[#test_func]] "test"
+; CHECK: OpName %[[#f2_decl:]] "BuiltInGlobalOffset"
+; CHECK: OpDecorate %[[#f2_decl]] LinkageAttributes "BuiltInGlobalOffset" Import
+; CHECK: %[[#int_ty:]] = OpTypeInt 32 0
+; CHECK: %[[#iptr_ty:]] = OpTypePointer CrossWorkgroup  %[[#int_ty]]
+; CHECK: %[[#void_ty:]] = OpTypeVoid
+; CHECK: %[[#func_ty:]] = OpTypeFunction %[[#void_ty]] %[[#iptr_ty]]
+; CHECK: %[[#int64_ty:]] = OpTypeInt 64 0
+; CHECK: %[[#vec_ty:]] = OpTypeVector %[[#int64_ty]] 3
+; CHECK: %[[#func2_ty:]] = OpTypeFunction %[[#vec_ty]]
+;; TODO: add 64-bit constant defs
+; CHECK: %[[#f2_decl]] = OpFunction %[[#vec_ty]] Pure %[[#func2_ty]]
+; CHECK: OpFunctionEnd
+;; Check that the function register name does not match other registers
+; CHECK-NOT: %[[#int_ty]] = OpFunction
+; CHECK-NOT: %[[#iptr_ty]] = OpFunction
+; CHECK-NOT: %[[#void_ty]] = OpFunction
+; CHECK-NOT: %[[#func_ty]] = OpFunction
+; CHECK-NOT: %[[#int64_ty]] = OpFunction
+; CHECK-NOT: %[[#vec_ty]] = OpFunction
+; CHECK-NOT: %[[#func2_ty]] = OpFunction
+; CHECK-NOT: %[[#f2_decl]] = OpFunction
+; CHECK: %[[#outOffsets]] = OpFunctionParameter %[[#iptr_ty]]
+
+define spir_kernel void @test(i32 addrspace(1)* %outOffsets) {
+entry:
+  %0 = call spir_func <3 x i64> @BuiltInGlobalOffset() #1
+  %call = extractelement <3 x i64> %0, i32 0
+  %conv = trunc i64 %call to i32
+; CHECK: %[[#i1:]] = OpInBoundsPtrAccessChain %[[#iptr_ty]] %[[#outOffsets]]
+; CHECK: OpStore %[[#i1:]] %[[#]] Aligned 4
+  %arrayidx = getelementptr inbounds i32, i32 addrspace(1)* %outOffsets, i64 0
+  store i32 %conv, i32 addrspace(1)* %arrayidx, align 4
+  %1 = call spir_func <3 x i64> @BuiltInGlobalOffset() #1
+  %call1 = extractelement <3 x i64> %1, i32 1
+  %conv2 = trunc i64 %call1 to i32
+; CHECK: %[[#i2:]] = OpInBoundsPtrAccessChain %[[#iptr_ty]] %[[#outOffsets]]
+; CHECK: OpStore %[[#i2:]] %[[#]] Aligned 4
+  %arrayidx3 = getelementptr inbounds i32, i32 addrspace(1)* %outOffsets, i64 1
+  store i32 %conv2, i32 addrspace(1)* %arrayidx3, align 4
+  %2 = call spir_func <3 x i64> @BuiltInGlobalOffset() #1
+  %call4 = extractelement <3 x i64> %2, i32 2
+  %conv5 = trunc i64 %call4 to i32
+; CHECK: %[[#i3:]] = OpInBoundsPtrAccessChain %[[#iptr_ty]] %[[#outOffsets]]
+; CHECK: OpStore %[[#i3:]] %[[#]] Aligned 4
+  %arrayidx6 = getelementptr inbounds i32, i32 addrspace(1)* %outOffsets, i64 2
+  store i32 %conv5, i32 addrspace(1)* %arrayidx6, align 4
+  ret void
+}
+
+declare spir_func <3 x i64> @BuiltInGlobalOffset() #1
+
+attributes #1 = { nounwind readnone }
diff --git a/llvm/test/CodeGen/SPIRV/opencl/basic/progvar_prog_scope_init.ll b/llvm/test/CodeGen/SPIRV/opencl/basic/progvar_prog_scope_init.ll
new file mode 100644
index 0000000000000..9d759a1cf47d0
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/opencl/basic/progvar_prog_scope_init.ll
@@ -0,0 +1,109 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+
+; CHECK:     OpEntryPoint Kernel %[[#f1:]] "writer"
+; CHECK:     OpEntryPoint Kernel %[[#f2:]] "reader"
+; CHECK-DAG: OpName %[[#a_var:]] "a_var"
+; CHECK-DAG: OpName %[[#p_var:]] "p_var"
+; CHECK-DAG: %[[#uchar:]] = OpTypeInt 8 0
+; CHECK-DAG: %[[#pt1:]] = OpTypePointer CrossWorkgroup %[[#uchar]]
+; CHECK-DAG: %[[#arr2:]] = OpTypeArray %[[#uchar]]
+; CHECK-DAG: %[[#pt2:]] = OpTypePointer CrossWorkgroup %[[#arr2]]
+; CHECK-DAG: %[[#pt3:]] = OpTypePointer CrossWorkgroup %[[#pt1]]
+; CHECK-DAG: %[[#a_var]] = OpVariable %[[#pt2]] CrossWorkgroup
+; CHECK-DAG: %[[#const:]] = OpSpecConstantOp %[[#pt1]] 70 %[[#a_var]]
+; CHECK-DAG: %[[#p_var]] = OpVariable %[[#pt3]] CrossWorkgroup %[[#const]]
+@var = addrspace(1) global i8 0, align 1
+@g_var = addrspace(1) global i8 1, align 1
+@a_var = addrspace(1) global [2 x i8] c"\01\01", align 1
+@p_var = addrspace(1) global i8 addrspace(1)* getelementptr inbounds ([2 x i8], [2 x i8] addrspace(1)* @a_var, i32 0, i64 1), align 8
+
+define spir_func zeroext i8 @from_buf(i8 zeroext %a) {
+entry:
+  %tobool = icmp ne i8 %a, 0
+  %i1promo = zext i1 %tobool to i8
+  ret i8 %i1promo
+}
+
+define spir_func zeroext i8 @to_buf(i8 zeroext %a) {
+entry:
+  %i1trunc = trunc i8 %a to i1
+  %frombool = select i1 %i1trunc, i8 1, i8 0
+  %0 = and i8 %frombool, 1
+  %tobool = icmp ne i8 %0, 0
+  %conv = select i1 %tobool, i8 1, i8 0
+  ret i8 %conv
+}
+
+define spir_kernel void @writer(i8 addrspace(1)* %src, i32 %idx) {
+entry:
+  %arrayidx = getelementptr inbounds i8, i8 addrspace(1)* %src, i64 0
+  %0 = load i8, i8 addrspace(1)* %arrayidx, align 1
+  %call = call spir_func zeroext i8 @from_buf(i8 zeroext %0)
+  %i1trunc = trunc i8 %call to i1
+  %frombool = select i1 %i1trunc, i8 1, i8 0
+  store i8 %frombool, i8 addrspace(1)* @var, align 1
+  %arrayidx1 = getelementptr inbounds i8, i8 addrspace(1)* %src, i64 1
+  %1 = load i8, i8 addrspace(1)* %arrayidx1, align 1
+  %call2 = call spir_func zeroext i8 @from_buf(i8 zeroext %1)
+  %i1trunc1 = trunc i8 %call2 to i1
+  %frombool3 = select i1 %i1trunc1, i8 1, i8 0
+  store i8 %frombool3, i8 addrspace(1)* @g_var, align 1
+  %arrayidx4 = getelementptr inbounds i8, i8 addrspace(1)* %src, i64 2
+  %2 = load i8, i8 addrspace(1)* %arrayidx4, align 1
+  %call5 = call spir_func zeroext i8 @from_buf(i8 zeroext %2)
+  %i1trunc2 = trunc i8 %call5 to i1
+  %frombool6 = select i1 %i1trunc2, i8 1, i8 0
+  %3 = getelementptr inbounds [2 x i8], [2 x i8] addrspace(1)* @a_var, i64 0, i64 0
+  store i8 %frombool6, i8 addrspace(1)* %3, align 1
+  %arrayidx7 = getelementptr inbounds i8, i8 addrspace(1)* %src, i64 3
+  %4 = load i8, i8 addrspace(1)* %arrayidx7, align 1
+  %call8 = call spir_func zeroext i8 @from_buf(i8 zeroext %4)
+  %i1trunc3 = trunc i8 %call8 to i1
+  %frombool9 = select i1 %i1trunc3, i8 1, i8 0
+  %5 = getelementptr inbounds [2 x i8], [2 x i8] addrspace(1)* @a_var, i64 0, i64 1
+  store i8 %frombool9, i8 addrspace(1)* %5, align 1
+  %idx.ext = zext i32 %idx to i64
+  %add.ptr = getelementptr inbounds i8, i8 addrspace(1)* %3, i64 %idx.ext
+  store i8 addrspace(1)* %add.ptr, i8 addrspace(1)* addrspace(1)* @p_var, align 8
+  ret void
+}
+
+define spir_kernel void @reader(i8 addrspace(1)* %dest, i8 zeroext %ptr_write_val) {
+entry:
+  %call = call spir_func zeroext i8 @from_buf(i8 zeroext %ptr_write_val)
+  %i1trunc = trunc i8 %call to i1
+  %0 = load i8 addrspace(1)*, i8 addrspace(1)* addrspace(1)* @p_var, align 8
+  %frombool = select i1 %i1trunc, i8 1, i8 0
+  store volatile i8 %frombool, i8 addrspace(1)* %0, align 1
+  %1 = load i8, i8 addrspace(1)* @var, align 1
+  %2 = and i8 %1, 1
+  %tobool = icmp ne i8 %2, 0
+  %i1promo = zext i1 %tobool to i8
+  %call1 = call spir_func zeroext i8 @to_buf(i8 zeroext %i1promo)
+  %arrayidx = getelementptr inbounds i8, i8 addrspace(1)* %dest, i64 0
+  store i8 %call1, i8 addrspace(1)* %arrayidx, align 1
+  %3 = load i8, i8 addrspace(1)* @g_var, align 1
+  %4 = and i8 %3, 1
+  %tobool2 = icmp ne i8 %4, 0
+  %i1promo1 = zext i1 %tobool2 to i8
+  %call3 = call spir_func zeroext i8 @to_buf(i8 zeroext %i1promo1)
+  %arrayidx4 = getelementptr inbounds i8, i8 addrspace(1)* %dest, i64 1
+  store i8 %call3, i8 addrspace(1)* %arrayidx4, align 1
+  %5 = getelementptr inbounds [2 x i8], [2 x i8] addrspace(1)* @a_var, i64 0, i64 0
+  %6 = load i8, i8 addrspace(1)* %5, align 1
+  %7 = and i8 %6, 1
+  %tobool5 = icmp ne i8 %7, 0
+  %i1promo2 = zext i1 %tobool5 to i8
+  %call6 = call spir_func zeroext i8 @to_buf(i8 zeroext %i1promo2)
+  %arrayidx7 = getelementptr inbounds i8, i8 addrspace(1)* %dest, i64 2
+  store i8 %call6, i8 addrspace(1)* %arrayidx7, align 1
+  %8 = getelementptr inbounds [2 x i8], [2 x i8] addrspace(1)* @a_var, i64 0, i64 1
+  %9 = load i8, i8 addrspace(1)* %8, align 1
+  %10 = and i8 %9, 1
+  %tobool8 = icmp ne i8 %10, 0
+  %i1promo3 = zext i1 %tobool8 to i8
+  %call9 = call spir_func zeroext i8 @to_buf(i8 zeroext %i1promo3)
+  %arrayidx10 = getelementptr inbounds i8, i8 addrspace(1)* %dest, i64 3
+  store i8 %call9, i8 addrspace(1)* %arrayidx10, align 1
+  ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/opencl/basic/progvar_prog_scope_uninit.ll b/llvm/test/CodeGen/SPIRV/opencl/basic/progvar_prog_scope_uninit.ll
new file mode 100644
index 0000000000000..fe02ba650aaa0
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/opencl/basic/progvar_prog_scope_uninit.ll
@@ -0,0 +1,152 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+
+; CHECK:     OpEntryPoint Kernel %[[#f1:]] "global_check" %[[#var0:]] %[[#var1:]] %[[#var2:]] %[[#var3:]]
+; CHECK:     OpEntryPoint Kernel %[[#f2:]] "writer" %[[#var0:]] %[[#var1:]] %[[#var2:]] %[[#var3:]]
+; CHECK:     OpEntryPoint Kernel %[[#f3:]] "reader" %[[#var0:]] %[[#var1:]] %[[#var2:]] %[[#var3:]]
+; CHECK-DAG: OpName %[[#var0]]
+; CHECK-DAG: OpName %[[#var1]]
+; CHECK-DAG: OpName %[[#var2]]
+; CHECK-DAG: OpName %[[#var3]]
+@var = addrspace(1) global <2 x i8> zeroinitializer, align 2
+@g_var = addrspace(1) global <2 x i8> zeroinitializer, align 2
+@a_var = addrspace(1) global [2 x <2 x i8>] zeroinitializer, align 2
+@p_var = addrspace(1) global <2 x i8> addrspace(1)* null, align 8
+
+define spir_func <2 x i8> @from_buf(<2 x i8> %a) {
+entry:
+  ret <2 x i8> %a
+}
+
+define spir_func <2 x i8> @to_buf(<2 x i8> %a) {
+entry:
+  ret <2 x i8> %a
+}
+
+define spir_kernel void @global_check(i32 addrspace(1)* %out) {
+entry:
+  %0 = load <2 x i8>, <2 x i8> addrspace(1)* @var, align 2
+  %cmp = icmp eq <2 x i8> %0, zeroinitializer
+  %sext = select <2 x i1> %cmp, <2 x i8> <i8 -1, i8 -1>, <2 x i8> zeroinitializer
+  %cast = icmp slt <2 x i8> %sext, zeroinitializer
+  %i1promo = zext <2 x i1> %cast to <2 x i8>
+  %call1 = call spir_func i1 @OpAll_v2i8(<2 x i8> %i1promo)
+  %call = select i1 %call1, i32 1, i32 0
+  %1 = and i8 1, 1
+  %tobool = icmp ne i8 %1, 0
+  %conv = select i1 %tobool, i32 1, i32 0
+  %and = and i32 %conv, %call
+  %tobool1 = icmp ne i32 %and, 0
+  %frombool = select i1 %tobool1, i8 1, i8 0
+  %2 = load <2 x i8>, <2 x i8> addrspace(1)* @g_var, align 2
+  %cmp2 = icmp eq <2 x i8> %2, zeroinitializer
+  %sext3 = select <2 x i1> %cmp2, <2 x i8> <i8 -1, i8 -1>, <2 x i8> zeroinitializer
+  %cast2 = icmp slt <2 x i8> %sext3, zeroinitializer
+  %i1promo1 = zext <2 x i1> %cast2 to <2 x i8>
+  %call43 = call spir_func i1 @OpAll_v2i8(<2 x i8> %i1promo1)
+  %call4 = select i1 %call43, i32 1, i32 0
+  %3 = and i8 %frombool, 1
+  %tobool5 = icmp ne i8 %3, 0
+  %conv6 = select i1 %tobool5, i32 1, i32 0
+  %and7 = and i32 %conv6, %call4
+  %tobool8 = icmp ne i32 %and7, 0
+  %frombool9 = select i1 %tobool8, i8 1, i8 0
+  %4 = getelementptr inbounds [2 x <2 x i8>], [2 x <2 x i8>] addrspace(1)* @a_var, i64 0, i64 0
+  %5 = load <2 x i8>, <2 x i8> addrspace(1)* %4, align 2
+  %cmp10 = icmp eq <2 x i8> %5, zeroinitializer
+  %sext11 = select <2 x i1> %cmp10, <2 x i8> <i8 -1, i8 -1>, <2 x i8> zeroinitializer
+  %cast4 = icmp slt <2 x i8> %sext11, zeroinitializer
+  %i1promo2 = zext <2 x i1> %cast4 to <2 x i8>
+  %call125 = call spir_func i1 @OpAll_v2i8(<2 x i8> %i1promo2)
+  %call12 = select i1 %call125, i32 1, i32 0
+  %6 = and i8 %frombool9, 1
+  %tobool13 = icmp ne i8 %6, 0
+  %conv14 = select i1 %tobool13, i32 1, i32 0
+  %and15 = and i32 %conv14, %call12
+  %tobool16 = icmp ne i32 %and15, 0
+  %frombool17 = select i1 %tobool16, i8 1, i8 0
+  %7 = getelementptr inbounds [2 x <2 x i8>], [2 x <2 x i8>] addrspace(1)* @a_var, i64 0, i64 1
+  %8 = load <2 x i8>, <2 x i8> addrspace(1)* %7, align 2
+  %cmp18 = icmp eq <2 x i8> %8, zeroinitializer
+  %sext19 = select <2 x i1> %cmp18, <2 x i8> <i8 -1, i8 -1>, <2 x i8> zeroinitializer
+  %cast6 = icmp slt <2 x i8> %sext19, zeroinitializer
+  %i1promo3 = zext <2 x i1> %cast6 to <2 x i8>
+  %call207 = call spir_func i1 @OpAll_v2i8(<2 x i8> %i1promo3)
+  %call20 = select i1 %call207, i32 1, i32 0
+  %9 = and i8 %frombool17, 1
+  %tobool21 = icmp ne i8 %9, 0
+  %conv22 = select i1 %tobool21, i32 1, i32 0
+  %and23 = and i32 %conv22, %call20
+  %tobool24 = icmp ne i32 %and23, 0
+  %frombool25 = select i1 %tobool24, i8 1, i8 0
+  %10 = load <2 x i8> addrspace(1)*, <2 x i8> addrspace(1)* addrspace(1)* @p_var, align 8
+  %11 = ptrtoint <2 x i8> addrspace(1)* %10 to i64
+  %12 = ptrtoint <2 x i8> addrspace(1)* null to i64
+  %cmp26 = icmp eq i64 %11, %12
+  %conv27 = select i1 %cmp26, i32 1, i32 0
+  %13 = and i8 %frombool25, 1
+  %tobool28 = icmp ne i8 %13, 0
+  %conv29 = select i1 %tobool28, i32 1, i32 0
+  %and30 = and i32 %conv29, %conv27
+  %tobool31 = icmp ne i32 %and30, 0
+  %frombool32 = select i1 %tobool31, i8 1, i8 0
+  %14 = and i8 %frombool32, 1
+  %tobool33 = icmp ne i8 %14, 0
+  %15 = select i1 %tobool33, i64 1, i64 0
+  %cond = select i1 %tobool33, i32 1, i32 0
+  store i32 %cond, i32 addrspace(1)* %out, align 4
+  ret void
+}
+
+declare spir_func i1 @OpAll_v2i8(<2 x i8>)
+
+define spir_kernel void @writer(<2 x i8> addrspace(1)* %src, i32 %idx) {
+entry:
+  %arrayidx = getelementptr inbounds <2 x i8>, <2 x i8> addrspace(1)* %src, i64 0
+  %0 = load <2 x i8>, <2 x i8> addrspace(1)* %arrayidx, align 2
+  %call = call spir_func <2 x i8> @from_buf(<2 x i8> %0)
+  store <2 x i8> %call, <2 x i8> addrspace(1)* @var, align 2
+  %arrayidx1 = getelementptr inbounds <2 x i8>, <2 x i8> addrspace(1)* %src, i64 1
+  %1 = load <2 x i8>, <2 x i8> addrspace(1)* %arrayidx1, align 2
+  %call2 = call spir_func <2 x i8> @from_buf(<2 x i8> %1)
+  store <2 x i8> %call2, <2 x i8> addrspace(1)* @g_var, align 2
+  %arrayidx3 = getelementptr inbounds <2 x i8>, <2 x i8> addrspace(1)* %src, i64 2
+  %2 = load <2 x i8>, <2 x i8> addrspace(1)* %arrayidx3, align 2
+  %call4 = call spir_func <2 x i8> @from_buf(<2 x i8> %2)
+  %3 = getelementptr inbounds [2 x <2 x i8>], [2 x <2 x i8>] addrspace(1)* @a_var, i64 0, i64 0
+  store <2 x i8> %call4, <2 x i8> addrspace(1)* %3, align 2
+  %arrayidx5 = getelementptr inbounds <2 x i8>, <2 x i8> addrspace(1)* %src, i64 3
+  %4 = load <2 x i8>, <2 x i8> addrspace(1)* %arrayidx5, align 2
+  %call6 = call spir_func <2 x i8> @from_buf(<2 x i8> %4)
+  %5 = getelementptr inbounds [2 x <2 x i8>], [2 x <2 x i8>] addrspace(1)* @a_var, i64 0, i64 1
+  store <2 x i8> %call6, <2 x i8> addrspace(1)* %5, align 2
+  %idx.ext = zext i32 %idx to i64
+  %add.ptr = getelementptr inbounds <2 x i8>, <2 x i8> addrspace(1)* %3, i64 %idx.ext
+  store <2 x i8> addrspace(1)* %add.ptr, <2 x i8> addrspace(1)* addrspace(1)* @p_var, align 8
+  ret void
+}
+
+define spir_kernel void @reader(<2 x i8> addrspace(1)* %dest, <2 x i8> %ptr_write_val) {
+entry:
+  %call = call spir_func <2 x i8> @from_buf(<2 x i8> %ptr_write_val)
+  %0 = load <2 x i8> addrspace(1)*, <2 x i8> addrspace(1)* addrspace(1)* @p_var, align 8
+  store <2 x i8> %call, <2 x i8> addrspace(1)* %0, align 2
+  %1 = load <2 x i8>, <2 x i8> addrspace(1)* @var, align 2
+  %call1 = call spir_func <2 x i8> @to_buf(<2 x i8> %1)
+  %arrayidx = getelementptr inbounds <2 x i8>, <2 x i8> addrspace(1)* %dest, i64 0
+  store <2 x i8> %call1, <2 x i8> addrspace(1)* %arrayidx, align 2
+  %2 = load <2 x i8>, <2 x i8> addrspace(1)* @g_var, align 2
+  %call2 = call spir_func <2 x i8> @to_buf(<2 x i8> %2)
+  %arrayidx3 = getelementptr inbounds <2 x i8>, <2 x i8> addrspace(1)* %dest, i64 1
+  store <2 x i8> %call2, <2 x i8> addrspace(1)* %arrayidx3, align 2
+  %3 = getelementptr inbounds [2 x <2 x i8>], [2 x <2 x i8>] addrspace(1)* @a_var, i64 0, i64 0
+  %4 = load <2 x i8>, <2 x i8> addrspace(1)* %3, align 2
+  %call4 = call spir_func <2 x i8> @to_buf(<2 x i8> %4)
+  %arrayidx5 = getelementptr inbounds <2 x i8>, <2 x i8> addrspace(1)* %dest, i64 2
+  store <2 x i8> %call4, <2 x i8> addrspace(1)* %arrayidx5, align 2
+  %5 = getelementptr inbounds [2 x <2 x i8>], [2 x <2 x i8>] addrspace(1)* @a_var, i64 0, i64 1
+  %6 = load <2 x i8>, <2 x i8> addrspace(1)* %5, align 2
+  %call6 = call spir_func <2 x i8> @to_buf(<2 x i8> %6)
+  %arrayidx7 = getelementptr inbounds <2 x i8>, <2 x i8> addrspace(1)* %dest, i64 3
+  store <2 x i8> %call6, <2 x i8> addrspace(1)* %arrayidx7, align 2
+  ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/preprocess-metadata.ll b/llvm/test/CodeGen/SPIRV/preprocess-metadata.ll
new file mode 100644
index 0000000000000..24e0851eb3833
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/preprocess-metadata.ll
@@ -0,0 +1,27 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+;; The purpose of this test is to check that some of OpenCL metadata are consumed
+;; even if 'opencl.ocl.version' metadata is missed (i.e. LLVM IR was produced not
+;; from OpenCL, but, for example from SYCL)
+
+; CHECK-SPIRV: OpEntryPoint Kernel %[[#TEST1:]] "test1"
+; CHECK-SPIRV: OpEntryPoint Kernel %[[#TEST2:]] "test2"
+; CHECK-SPIRV: OpExecutionMode %[[#TEST1]] LocalSize 1 2 3
+; CHECK-SPIRV: OpExecutionMode %[[#TEST1]] VecTypeHint 6
+; CHECK-SPIRV: OpExecutionMode %[[#TEST2]] LocalSizeHint 3 2 1
+; CHECK-SPIRV: OpExecutionMode %[[#TEST2]] SubgroupSize 8
+
+define spir_kernel void @test1() !reqd_work_group_size !1 !vec_type_hint !2 {
+entry:
+  ret void
+}
+
+define spir_kernel void @test2() !work_group_size_hint !3 !intel_reqd_sub_group_size !4 {
+entry:
+  ret void
+}
+
+!1 = !{i32 1, i32 2, i32 3}
+!2 = !{double undef, i32 1}
+!3 = !{i32 3, i32 2, i32 1}
+!4 = !{i32 8}
diff --git a/llvm/test/CodeGen/SPIRV/pstruct.ll b/llvm/test/CodeGen/SPIRV/pstruct.ll
new file mode 100644
index 0000000000000..9ae80ab2ab54f
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/pstruct.ll
@@ -0,0 +1,121 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+%struct.ST = type { i32, i32, i32 }
+
+; CHECK-SPIRV: OpName %[[#struct:]] "struct.ST"
+; CHECK-SPIRV: %[[#int:]] = OpTypeInt 32 0
+; CHECK-SPIRV: %[[#intP:]] = OpTypePointer Function %[[#int]]
+; CHECK-SPIRV: %[[#struct]] = OpTypeStruct %[[#int]] %[[#int]] %[[#int]]
+; CHECK-SPIRV: %[[#structP:]] = OpTypePointer Function %[[#struct]]
+; CHECK-SPIRV: %[[#structPP:]] = OpTypePointer Function %[[#structP]]
+; CHECK-SPIRV: %[[#zero:]] = OpConstant %[[#int]] 0
+; CHECK-SPIRV: %[[#one:]] = OpConstant %[[#int]] 1
+; CHECK-SPIRV: %[[#two:]] = OpConstant %[[#int]] 2
+
+define dso_local spir_func i32 @cmp_func(i8* %p1, i8* %p2) {
+entry:
+  %retval = alloca i32, align 4
+  %p1.addr = alloca i8*, align 8
+  %p2.addr = alloca i8*, align 8
+; CHECK-SPIRV: %[[#s1:]] = OpVariable %[[#structPP]]
+; CHECK-SPIRV: %[[#s2:]] = OpVariable %[[#structPP]]
+  %s1 = alloca %struct.ST*, align 8
+  %s2 = alloca %struct.ST*, align 8
+  store i8* %p1, i8** %p1.addr, align 8
+  store i8* %p2, i8** %p2.addr, align 8
+  %0 = load i8*, i8** %p1.addr, align 8
+; CHECK-SPIRV: %[[#t1:]] = OpBitcast %[[#structP]]
+; CHECK-SPIRV: OpStore %[[#s1]] %[[#t1]]
+  %1 = bitcast i8* %0 to %struct.ST*
+  store %struct.ST* %1, %struct.ST** %s1, align 8
+  %2 = load i8*, i8** %p2.addr, align 8
+; CHECK-SPIRV: %[[#t2:]] = OpBitcast %[[#structP]]
+; CHECK-SPIRV: OpStore %[[#s2]] %[[#t2]]
+  %3 = bitcast i8* %2 to %struct.ST*
+  store %struct.ST* %3, %struct.ST** %s2, align 8
+; CHECK-SPIRV: %[[#t3:]] = OpLoad %[[#structP]] %[[#s1]]
+; CHECK-SPIRV: %[[#a1:]] = OpInBoundsPtrAccessChain %[[#intP]] %[[#t3]] %[[#zero]] %[[#zero]]
+; CHECK-SPIRV: %[[#]] = OpLoad %[[#int]] %[[#a1]]
+  %4 = load %struct.ST*, %struct.ST** %s1, align 8
+  %a = getelementptr inbounds %struct.ST, %struct.ST* %4, i32 0, i32 0
+  %5 = load i32, i32* %a, align 4
+; CHECK-SPIRV: %[[#t4:]] = OpLoad %[[#structP]] %[[#s2]]
+; CHECK-SPIRV: %[[#a2:]] = OpInBoundsPtrAccessChain %[[#intP]] %[[#t4]] %[[#zero]] %[[#zero]]
+; CHECK-SPIRV: %[[#]] = OpLoad %[[#int]] %[[#a2]]
+  %6 = load %struct.ST*, %struct.ST** %s2, align 8
+  %a1 = getelementptr inbounds %struct.ST, %struct.ST* %6, i32 0, i32 0
+  %7 = load i32, i32* %a1, align 4
+  %cmp = icmp ne i32 %5, %7
+  br i1 %cmp, label %if.then, label %if.end
+
+if.then:                                          ; preds = %entry
+; CHECK-SPIRV: %[[#t5:]] = OpLoad %[[#structP]] %[[#s1]]
+; CHECK-SPIRV: %[[#a_1:]] = OpInBoundsPtrAccessChain %[[#intP]] %[[#t5]] %[[#zero]] %[[#zero]]
+; CHECK-SPIRV: %[[#]] = OpLoad %[[#int]] %[[#a_1]]
+  %8 = load %struct.ST*, %struct.ST** %s1, align 8
+  %a2 = getelementptr inbounds %struct.ST, %struct.ST* %8, i32 0, i32 0
+  %9 = load i32, i32* %a2, align 4
+; CHECK-SPIRV: %[[#t6:]] = OpLoad %[[#structP]] %[[#s2]]
+; CHECK-SPIRV: %[[#a_2:]] = OpInBoundsPtrAccessChain %[[#intP]] %[[#t6]] %[[#zero]] %[[#zero]]
+; CHECK-SPIRV: %[[#]] = OpLoad %[[#int]] %[[#a_2]]
+  %10 = load %struct.ST*, %struct.ST** %s2, align 8
+  %a3 = getelementptr inbounds %struct.ST, %struct.ST* %10, i32 0, i32 0
+  %11 = load i32, i32* %a3, align 4
+  %sub = sub nsw i32 %9, %11
+  store i32 %sub, i32* %retval, align 4
+  br label %return
+
+if.end:                                           ; preds = %entry
+; CHECK-SPIRV: %[[#t7:]] = OpLoad %[[#structP]] %[[#s1]]
+; CHECK-SPIRV: %[[#b1:]] = OpInBoundsPtrAccessChain %[[#intP]] %[[#t7]] %[[#zero]] %[[#one]]
+; CHECK-SPIRV: %[[#]] = OpLoad %[[#int]] %[[#b1]]
+  %12 = load %struct.ST*, %struct.ST** %s1, align 8
+  %b = getelementptr inbounds %struct.ST, %struct.ST* %12, i32 0, i32 1
+  %13 = load i32, i32* %b, align 4
+; CHECK-SPIRV: %[[#t8:]] = OpLoad %[[#structP]] %[[#s2]]
+; CHECK-SPIRV: %[[#b2:]] = OpInBoundsPtrAccessChain %[[#intP]] %[[#t8]] %[[#zero]] %[[#one]]
+; CHECK-SPIRV: %[[#]] = OpLoad %[[#int]] %[[#b2]]
+  %14 = load %struct.ST*, %struct.ST** %s2, align 8
+  %b4 = getelementptr inbounds %struct.ST, %struct.ST* %14, i32 0, i32 1
+  %15 = load i32, i32* %b4, align 4
+  %cmp5 = icmp ne i32 %13, %15
+  br i1 %cmp5, label %if.then6, label %if.end10
+
+if.then6:                                         ; preds = %if.end
+; CHECK-SPIRV: %[[#t9:]] = OpLoad %[[#structP]] %[[#s1]]
+; CHECK-SPIRV: %[[#b_1:]] = OpInBoundsPtrAccessChain %[[#intP]] %[[#t9]] %[[#zero]] %[[#one]]
+; CHECK-SPIRV: %[[#]] = OpLoad %[[#int]] %[[#b_1]]
+  %16 = load %struct.ST*, %struct.ST** %s1, align 8
+  %b7 = getelementptr inbounds %struct.ST, %struct.ST* %16, i32 0, i32 1
+  %17 = load i32, i32* %b7, align 4
+; CHECK-SPIRV: %[[#t10:]] = OpLoad %[[#structP]] %[[#s2]]
+; CHECK-SPIRV: %[[#b_2:]] = OpInBoundsPtrAccessChain %[[#intP]] %[[#t10]] %[[#zero]] %[[#one]]
+; CHECK-SPIRV: %[[#]] = OpLoad %[[#int]] %[[#b_2]]
+  %18 = load %struct.ST*, %struct.ST** %s2, align 8
+  %b8 = getelementptr inbounds %struct.ST, %struct.ST* %18, i32 0, i32 1
+  %19 = load i32, i32* %b8, align 4
+  %sub9 = sub nsw i32 %17, %19
+  store i32 %sub9, i32* %retval, align 4
+  br label %return
+
+if.end10:                                         ; preds = %if.end
+; CHECK-SPIRV: %[[#t11:]] = OpLoad %[[#structP]] %[[#s1]]
+; CHECK-SPIRV: %[[#c1:]] = OpInBoundsPtrAccessChain %[[#intP]] %[[#t11]] %[[#zero]] %[[#two]]
+; CHECK-SPIRV: %[[#]] = OpLoad %[[#int]] %[[#c1]]
+  %20 = load %struct.ST*, %struct.ST** %s1, align 8
+  %c = getelementptr inbounds %struct.ST, %struct.ST* %20, i32 0, i32 2
+  %21 = load i32, i32* %c, align 4
+; CHECK-SPIRV: %[[#t12:]] = OpLoad %[[#structP]] %[[#s2]]
+; CHECK-SPIRV: %[[#c2:]] = OpInBoundsPtrAccessChain %[[#intP]] %[[#t12]] %[[#zero]] %[[#two]]
+; CHECK-SPIRV: %[[#]] = OpLoad %[[#int]] %[[#c2]]
+  %22 = load %struct.ST*, %struct.ST** %s2, align 8
+  %c11 = getelementptr inbounds %struct.ST, %struct.ST* %22, i32 0, i32 2
+  %23 = load i32, i32* %c11, align 4
+  %sub12 = sub nsw i32 %21, %23
+  store i32 %sub12, i32* %retval, align 4
+  br label %return
+
+return:                                           ; preds = %if.end10, %if.then6, %if.then
+  %24 = load i32, i32* %retval, align 4
+  ret i32 %24
+}
diff --git a/llvm/test/CodeGen/SPIRV/sitofp-with-bool.ll b/llvm/test/CodeGen/SPIRV/sitofp-with-bool.ll
new file mode 100644
index 0000000000000..4dda0aadced49
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/sitofp-with-bool.ll
@@ -0,0 +1,22 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+
+; CHECK: %[[#int_32:]] = OpTypeInt 32 0
+; CHECK: %[[#bool:]] = OpTypeBool
+; CHECK: %[[#zero:]] = OpConstant %[[#int_32]] 0
+; CHECK: %[[#one:]] = OpConstant %[[#int_32]] 1
+
+; CHECK: OpFunction
+; CHECK: %[[#A:]] = OpFunctionParameter %[[#]]
+; CHECK: %[[#B:]] = OpFunctionParameter %[[#]]
+; CHECK: %[[#cmp_res:]] = OpSGreaterThan %[[#bool]] %[[#B]] %[[#zero]]
+; CHECK: %[[#select_res:]] = OpSelect %[[#int_32]] %[[#cmp_res]] %[[#one]] %[[#zero]]
+; CHECK: %[[#stof_res:]] = OpConvertSToF %[[#]] %[[#select_res]]
+; CHECK: OpStore %[[#A]] %[[#stof_res]]
+
+define dso_local spir_kernel void @K(float addrspace(1)* nocapture %A, i32 %B) local_unnamed_addr {
+entry:
+  %cmp = icmp sgt i32 %B, 0
+  %conv = sitofp i1 %cmp to float
+  store float %conv, float addrspace(1)* %A, align 4
+  ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/spirv_param_decorations.ll b/llvm/test/CodeGen/SPIRV/spirv_param_decorations.ll
new file mode 100644
index 0000000000000..dbd8336d45809
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/spirv_param_decorations.ll
@@ -0,0 +1,20 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+define spir_kernel void @k(float %a, float %b, float %c) !spirv.ParameterDecorations !14 {
+entry:
+  ret void
+}
+
+; CHECK-SPIRV: OpDecorate %[[#PId1:]] Restrict
+; CHECK-SPIRV: OpDecorate %[[#PId1]] FPRoundingMode RTP
+; CHECK-SPIRV: OpDecorate %[[#PId2:]] Volatile
+; CHECK-SPIRV: %[[#PId1]] = OpFunctionParameter %[[#]]
+; CHECK-SPIRV: %[[#PId2]] = OpFunctionParameter %[[#]]
+
+!8 = !{i32 19}
+!9 = !{i32 39, i32 2}
+!10 = !{i32 21}
+!11 = !{!8, !9}
+!12 = !{}
+!13 = !{!10}
+!14 = !{!11, !12, !13}
diff --git a/llvm/test/CodeGen/SPIRV/spirv_param_decorations_quals.ll b/llvm/test/CodeGen/SPIRV/spirv_param_decorations_quals.ll
new file mode 100644
index 0000000000000..260394b658348
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/spirv_param_decorations_quals.ll
@@ -0,0 +1,15 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+define spir_kernel void @k(i32 addrspace(1)* %a) !kernel_arg_type_qual !7 !spirv.ParameterDecorations !10 {
+entry:
+  ret void
+}
+
+; CHECK-SPIRV: OpDecorate %[[#PId:]] Volatile
+; CHECK-SPIRV: OpDecorate %[[#PId]] FuncParamAttr NoAlias
+; CHECK-SPIRV: %[[#PId]] = OpFunctionParameter %[[#]]
+
+!7 = !{!"volatile"}
+!8 = !{i32 38, i32 4} ; FuncParamAttr NoAlias
+!9 = !{!8}
+!10 = !{!9}
diff --git a/llvm/test/CodeGen/SPIRV/store.ll b/llvm/test/CodeGen/SPIRV/store.ll
new file mode 100644
index 0000000000000..386a605dc4aa5
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/store.ll
@@ -0,0 +1,12 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
+
+; CHECK: "foo"
+define spir_kernel void @foo(i32 addrspace(1)* %a) {
+entry:
+  %a.addr = alloca i32 addrspace(1)*, align 4
+  store i32 addrspace(1)* %a, i32 addrspace(1)** %a.addr, align 4
+  %0 = load i32 addrspace(1)*, i32 addrspace(1)** %a.addr, align 4
+; CHECK: OpStore %[[#]] %[[#]] Aligned 4
+  store i32 0, i32 addrspace(1)* %0, align 4
+  ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/struct.ll b/llvm/test/CodeGen/SPIRV/struct.ll
new file mode 100644
index 0000000000000..56b6c9f9ee826
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/struct.ll
@@ -0,0 +1,46 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+%struct.ST = type { i32, i32, i32 }
+
+; CHECK-SPIRV: OpName %[[#struct:]] "struct.ST"
+; CHECK-SPIRV: %[[#int:]] = OpTypeInt 32 0
+; CHECK-SPIRV-DAG: %[[#struct]] = OpTypeStruct %[[#int]] %[[#int]] %[[#int]]
+; CHECK-SPIRV-DAG: %[[#structP:]] = OpTypePointer Function %[[#struct]]
+; CHECK-SPIRV-DAG: %[[#intP:]] = OpTypePointer Function %[[#int]]
+; CHECK-SPIRV: %[[#zero:]] = OpConstant %[[#int]] 0
+; CHECK-SPIRV: %[[#one:]] = OpConstant %[[#int]] 1
+; CHECK-SPIRV: %[[#two:]] = OpConstant %[[#int]] 2
+; CHECK-SPIRV: %[[#three:]] = OpConstant %[[#int]] 3
+
+define dso_local spir_func i32 @func() {
+entry:
+; CHECK-SPIRV: %[[#st:]] = OpVariable %[[#structP]]
+  %st = alloca %struct.ST, align 4
+; CHECK-SPIRV: %[[#a:]] = OpInBoundsPtrAccessChain %[[#intP]] %[[#st]] %[[#zero]] %[[#zero]]
+; CHECK-SPIRV: OpStore %[[#a]] %[[#one]]
+  %a = getelementptr inbounds %struct.ST, %struct.ST* %st, i32 0, i32 0
+  store i32 1, i32* %a, align 4
+; CHECK-SPIRV: %[[#b:]] = OpInBoundsPtrAccessChain %[[#intP]] %[[#st]] %[[#zero]] %[[#one]]
+; CHECK-SPIRV: OpStore %[[#b]] %[[#two]]
+  %b = getelementptr inbounds %struct.ST, %struct.ST* %st, i32 0, i32 1
+  store i32 2, i32* %b, align 4
+; CHECK-SPIRV: %[[#c:]] = OpInBoundsPtrAccessChain %[[#intP]] %[[#st]] %[[#zero]] %[[#two]]
+; CHECK-SPIRV: OpStore %[[#c]] %[[#three]]
+  %c = getelementptr inbounds %struct.ST, %struct.ST* %st, i32 0, i32 2
+  store i32 3, i32* %c, align 4
+; CHECK-SPIRV: %[[#a1:]] = OpInBoundsPtrAccessChain %[[#intP]] %[[#st]] %[[#zero]] %[[#zero]]
+; CHECK-SPIRV: %[[#]] = OpLoad %[[#int]] %[[#a1]]
+  %a1 = getelementptr inbounds %struct.ST, %struct.ST* %st, i32 0, i32 0
+  %0 = load i32, i32* %a1, align 4
+; CHECK-SPIRV: %[[#b1:]] = OpInBoundsPtrAccessChain %[[#intP]] %[[#st]] %[[#zero]] %[[#one]]
+; CHECK-SPIRV: %[[#]] = OpLoad %[[#int]] %[[#b1]]
+  %b2 = getelementptr inbounds %struct.ST, %struct.ST* %st, i32 0, i32 1
+  %1 = load i32, i32* %b2, align 4
+  %add = add nsw i32 %0, %1
+; CHECK-SPIRV: %[[#c1:]] = OpInBoundsPtrAccessChain %[[#intP]] %[[#st]] %[[#zero]] %[[#two]]
+; CHECK-SPIRV: %[[#]] = OpLoad %[[#int]] %[[#c1]]
+  %c3 = getelementptr inbounds %struct.ST, %struct.ST* %st, i32 0, i32 2
+  %2 = load i32, i32* %c3, align 4
+  %add4 = add nsw i32 %add, %2
+  ret i32 %add4
+}
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/ConvertPtr.ll b/llvm/test/CodeGen/SPIRV/transcoding/ConvertPtr.ll
new file mode 100644
index 0000000000000..34036951e31e0
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/ConvertPtr.ll
@@ -0,0 +1,30 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+;; kernel void testConvertPtrToU(global int *a, global unsigned long *res) {
+;;   res[0] = (unsigned long)&a[0];
+;; }
+
+; CHECK-SPIRV: OpConvertPtrToU
+
+define dso_local spir_kernel void @testConvertPtrToU(i32 addrspace(1)* noundef %a, i64 addrspace(1)* nocapture noundef writeonly %res) local_unnamed_addr {
+entry:
+  %0 = ptrtoint i32 addrspace(1)* %a to i32
+  %1 = zext i32 %0 to i64
+  store i64 %1, i64 addrspace(1)* %res, align 8
+  ret void
+}
+
+;; kernel void testConvertUToPtr(unsigned long a) {
+;;   global unsigned int *res = (global unsigned int *)a;
+;;   res[0] = 0;
+;; }
+
+; CHECK-SPIRV: OpConvertUToPtr
+
+define dso_local spir_kernel void @testConvertUToPtr(i64 noundef %a) local_unnamed_addr {
+entry:
+  %conv = trunc i64 %a to i32
+  %0 = inttoptr i32 %conv to i32 addrspace(1)*
+  store i32 0, i32 addrspace(1)* %0, align 4
+  ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/DecorationAlignment.ll b/llvm/test/CodeGen/SPIRV/transcoding/DecorationAlignment.ll
new file mode 100644
index 0000000000000..2e9b4a494c04d
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/DecorationAlignment.ll
@@ -0,0 +1,11 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: OpDecorate %[[#ALIGNMENT:]] Alignment 16
+; CHECK-SPIRV: %[[#ALIGNMENT]] = OpFunctionParameter %[[#]]
+
+%struct._ZTS6Struct.Struct = type { %struct._ZTS11floatStruct.floatStruct, %struct._ZTS11floatStruct.floatStruct }
+%struct._ZTS11floatStruct.floatStruct = type { float, float, float, float }
+
+define spir_func void @_ZN3FooC2Ev(%struct._ZTS6Struct.Struct addrspace(4)* align 16 %0) {
+  ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/DivRem.ll b/llvm/test/CodeGen/SPIRV/transcoding/DivRem.ll
new file mode 100644
index 0000000000000..2f423c2518e83
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/DivRem.ll
@@ -0,0 +1,91 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV-DAG: %[[#int:]] = OpTypeInt 32 0
+; CHECK-SPIRV-DAG: %[[#int2:]] = OpTypeVector %[[#int]] 2
+; CHECK-SPIRV-DAG: %[[#float:]] = OpTypeFloat 32
+; CHECK-SPIRV-DAG: %[[#float2:]] = OpTypeVector %[[#float]] 2
+
+; CHECK-SPIRV:      OpFunction
+; CHECK-SPIRV-NEXT: %[[#A:]] = OpFunctionParameter %[[#]]
+; CHECK-SPIRV-NEXT: %[[#B:]] = OpFunctionParameter %[[#]]
+; CHECK-SPIRV:      %[[#]] = OpSDiv %[[#int2]] %[[#A]] %[[#B]]
+; CHECK-SPIRV:      OpFunctionEnd
+
+;; kernel void testSDiv(int2 a, int2 b, global int2 *res) {
+;;   res[0] = a / b;
+;; }
+
+define dso_local spir_kernel void @testSDiv(<2 x i32> noundef %a, <2 x i32> noundef %b, <2 x i32> addrspace(1)* nocapture noundef writeonly %res) local_unnamed_addr {
+entry:
+  %div = sdiv <2 x i32> %a, %b
+  store <2 x i32> %div, <2 x i32> addrspace(1)* %res, align 8
+  ret void
+}
+
+; CHECK-SPIRV:      OpFunction
+; CHECK-SPIRV-NEXT: %[[#A:]] = OpFunctionParameter %[[#]]
+; CHECK-SPIRV-NEXT: %[[#B:]] = OpFunctionParameter %[[#]]
+; CHECK-SPIRV:      %[[#]] = OpUDiv %[[#int2]] %[[#A]] %[[#B]]
+; CHECK-SPIRV:      OpFunctionEnd
+
+;; kernel void testUDiv(uint2 a, uint2 b, global uint2 *res) {
+;;   res[0] = a / b;
+;; }
+
+define dso_local spir_kernel void @testUDiv(<2 x i32> noundef %a, <2 x i32> noundef %b, <2 x i32> addrspace(1)* nocapture noundef writeonly %res) local_unnamed_addr {
+entry:
+  %div = udiv <2 x i32> %a, %b
+  store <2 x i32> %div, <2 x i32> addrspace(1)* %res, align 8
+  ret void
+}
+
+; CHECK-SPIRV:      OpFunction
+; CHECK-SPIRV-NEXT: %[[#A:]] = OpFunctionParameter %[[#]]
+; CHECK-SPIRV-NEXT: %[[#B:]] = OpFunctionParameter %[[#]]
+; CHECK-SPIRV:      %[[#]] = OpFDiv %[[#float2]] %[[#A]] %[[#B]]
+; CHECK-SPIRV:      OpFunctionEnd
+
+;; kernel void testFDiv(float2 a, float2 b, global float2 *res) {
+;;   res[0] = a / b;
+;; }
+
+define dso_local spir_kernel void @testFDiv(<2 x float> noundef %a, <2 x float> noundef %b, <2 x float> addrspace(1)* nocapture noundef writeonly %res) local_unnamed_addr {
+entry:
+  %div = fdiv <2 x float> %a, %b
+  store <2 x float> %div, <2 x float> addrspace(1)* %res, align 8
+  ret void
+}
+
+; CHECK-SPIRV:      OpFunction
+; CHECK-SPIRV-NEXT: %[[#A:]] = OpFunctionParameter %[[#]]
+; CHECK-SPIRV-NEXT: %[[#B:]] = OpFunctionParameter %[[#]]
+; CHECK-SPIRV:      %[[#]] = OpSRem %[[#int2]] %[[#A]] %[[#B]]
+; CHECK-SPIRV:      OpFunctionEnd
+
+;; kernel void testSRem(int2 a, int2 b, global int2 *res) {
+;;   res[0] = a % b;
+;; }
+
+define dso_local spir_kernel void @testSRem(<2 x i32> noundef %a, <2 x i32> noundef %b, <2 x i32> addrspace(1)* nocapture noundef writeonly %res) local_unnamed_addr {
+entry:
+  %rem = srem <2 x i32> %a, %b
+  store <2 x i32> %rem, <2 x i32> addrspace(1)* %res, align 8
+  ret void
+}
+
+; CHECK-SPIRV:      OpFunction
+; CHECK-SPIRV-NEXT: %[[#A:]] = OpFunctionParameter %[[#]]
+; CHECK-SPIRV-NEXT: %[[#B:]] = OpFunctionParameter %[[#]]
+; CHECK-SPIRV:      %[[#]] = OpUMod %[[#int2]] %[[#A]] %[[#B]]
+; CHECK-SPIRV:      OpFunctionEnd
+
+;; kernel void testUMod(uint2 a, uint2 b, global uint2 *res) {
+;;   res[0] = a % b;
+;; }
+
+define dso_local spir_kernel void @testUMod(<2 x i32> noundef %a, <2 x i32> noundef %b, <2 x i32> addrspace(1)* nocapture noundef writeonly %res) local_unnamed_addr {
+entry:
+  %rem = urem <2 x i32> %a, %b
+  store <2 x i32> %rem, <2 x i32> addrspace(1)* %res, align 8
+  ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/ExecutionMode_SPIR_to_SPIRV.ll b/llvm/test/CodeGen/SPIRV/transcoding/ExecutionMode_SPIR_to_SPIRV.ll
new file mode 100644
index 0000000000000..6d6dd2481b17d
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/ExecutionMode_SPIR_to_SPIRV.ll
@@ -0,0 +1,11 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV-DAG: OpEntryPoint Kernel %[[#WORKER:]] "worker"
+; CHECK-SPIRV-DAG: OpExecutionMode %[[#WORKER]] LocalSizeHint 128 10 1
+
+define spir_kernel void @worker() local_unnamed_addr !work_group_size_hint !3 {
+entry:
+  ret void
+}
+
+!3 = !{i32 128, i32 10, i32 1}
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/GlobalFunAnnotate.ll b/llvm/test/CodeGen/SPIRV/transcoding/GlobalFunAnnotate.ll
new file mode 100644
index 0000000000000..2796dcbdca948
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/GlobalFunAnnotate.ll
@@ -0,0 +1,12 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-linux %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: OpDecorate %[[#]] UserSemantic "annotation_on_function"
+
+@.str = private unnamed_addr constant [23 x i8] c"annotation_on_function\00", section "llvm.metadata"
+@.str.1 = private unnamed_addr constant [6 x i8] c"an.cl\00", section "llvm.metadata"
+@llvm.global.annotations = appending global [1 x { i8*, i8*, i8*, i32, i8* }] [{ i8*, i8*, i8*, i32, i8* } { i8* bitcast (void ()* @foo to i8*), i8* getelementptr inbounds ([23 x i8], [23 x i8]* @.str, i32 0, i32 0), i8* getelementptr inbounds ([6 x i8], [6 x i8]* @.str.1, i32 0, i32 0), i32 2, i8* null }], section "llvm.metadata"
+
+define dso_local spir_func void @foo() {
+entry:
+  ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpConstantBool.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpConstantBool.ll
new file mode 100644
index 0000000000000..c0d61954de45b
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpConstantBool.ll
@@ -0,0 +1,25 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: OpConstantTrue
+; CHECK-SPIRV: OpConstantFalse
+
+define spir_func zeroext i1 @f() {
+entry:
+  ret i1 true
+}
+
+define spir_func zeroext i1 @f2() {
+entry:
+  ret i1 false
+}
+
+define spir_kernel void @test(i32 addrspace(1)* %i) {
+entry:
+  %i.addr = alloca i32 addrspace(1)*, align 4
+  store i32 addrspace(1)* %i, i32 addrspace(1)** %i.addr, align 4
+  %call = call spir_func zeroext i1 @f()
+  %conv = zext i1 %call to i32
+  %0 = load i32 addrspace(1)*, i32 addrspace(1)** %i.addr, align 4
+  store i32 %conv, i32 addrspace(1)* %0, align 4
+  ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpPhi_ArgumentsPlaceholders.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpPhi_ArgumentsPlaceholders.ll
new file mode 100644
index 0000000000000..a4f5530b6534f
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpPhi_ArgumentsPlaceholders.ll
@@ -0,0 +1,49 @@
+;; struct Node;
+;; typedef struct {
+;;     __global struct Node* pNext;
+;; } Node;
+;;
+;; __kernel void verify_linked_lists(__global Node* pNodes)
+;; {
+;;     __global Node *pNode = pNodes;
+;;
+;;     for(int j=0; j < 10; j++) {
+;;         pNode = pNode->pNext;
+;;     }
+;; }
+
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+%struct.Node = type { %struct.Node.0 addrspace(1)* }
+%struct.Node.0 = type opaque
+
+define spir_kernel void @verify_linked_lists(%struct.Node addrspace(1)* %pNodes) {
+entry:
+  br label %for.cond
+
+for.cond:                                         ; preds = %for.inc, %entry
+  %pNode.0 = phi %struct.Node addrspace(1)* [ %pNodes, %entry ], [ %1, %for.inc ]
+  %j.0 = phi i32 [ 0, %entry ], [ %inc, %for.inc ]
+;CHECK-SPIRV: %[[#]] = OpPhi %[[#]] %[[#]] %[[#]] %[[#BitcastResultId:]] %[[#]]
+;CHECK-SPIRV-NEXT: OpPhi
+
+  %cmp = icmp slt i32 %j.0, 10
+  br i1 %cmp, label %for.body, label %for.end
+
+for.body:                                         ; preds = %for.cond
+  %pNext = getelementptr inbounds %struct.Node, %struct.Node addrspace(1)* %pNode.0, i32 0, i32 0
+
+  %0 = load %struct.Node.0 addrspace(1)*, %struct.Node.0 addrspace(1)* addrspace(1)* %pNext, align 4
+  %1 = bitcast %struct.Node.0 addrspace(1)* %0 to %struct.Node addrspace(1)*
+;CHECK-SPIRV: %[[#LoadResultId:]] = OpLoad %[[#]]
+;CHECK-SPIRV: %[[#BitcastResultId]] = OpBitcast %[[#]] %[[#LoadResultId]]
+
+  br label %for.inc
+
+for.inc:                                          ; preds = %for.body
+  %inc = add nsw i32 %j.0, 1
+  br label %for.cond
+
+for.end:                                          ; preds = %for.cond
+  ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpSwitch32.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpSwitch32.ll
new file mode 100644
index 0000000000000..fdf4e4ddbefb7
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpSwitch32.ll
@@ -0,0 +1,55 @@
+;; __kernel void test_32(__global int* res)
+;; {
+;;     int tid = get_global_id(0);
+;;
+;;     switch(tid)
+;;     {
+;;     case 0:
+;;         res[tid] = 1;
+;;         break;
+;;     case 1:
+;;         res[tid] = 2;
+;;         break;
+;;     }
+;; }
+;; bash$ clang -cc1 -triple spir64-unknown-unknown -x cl -cl-std=CL2.0 -O0 -include opencl.h -emit-llvm OpSwitch.cl -o test_32.ll
+
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: OpSwitch %[[#]] %[[#]] 0 %[[#]] 1 %[[#]]
+
+define spir_kernel void @test_32(i32 addrspace(1)* %res) {
+entry:
+  %res.addr = alloca i32 addrspace(1)*, align 8
+  %tid = alloca i32, align 4
+  store i32 addrspace(1)* %res, i32 addrspace(1)** %res.addr, align 8
+  %call = call spir_func i64 @_Z13get_global_idj(i32 0)
+  %conv = trunc i64 %call to i32
+  store i32 %conv, i32* %tid, align 4
+  %0 = load i32, i32* %tid, align 4
+  switch i32 %0, label %sw.epilog [
+    i32 0, label %sw.bb
+    i32 1, label %sw.bb1
+  ]
+
+sw.bb:                                            ; preds = %entry
+  %1 = load i32, i32* %tid, align 4
+  %idxprom = sext i32 %1 to i64
+  %2 = load i32 addrspace(1)*, i32 addrspace(1)** %res.addr, align 8
+  %arrayidx = getelementptr inbounds i32, i32 addrspace(1)* %2, i64 %idxprom
+  store i32 1, i32 addrspace(1)* %arrayidx, align 4
+  br label %sw.epilog
+
+sw.bb1:                                           ; preds = %entry
+  %3 = load i32, i32* %tid, align 4
+  %idxprom2 = sext i32 %3 to i64
+  %4 = load i32 addrspace(1)*, i32 addrspace(1)** %res.addr, align 8
+  %arrayidx3 = getelementptr inbounds i32, i32 addrspace(1)* %4, i64 %idxprom2
+  store i32 2, i32 addrspace(1)* %arrayidx3, align 4
+  br label %sw.epilog
+
+sw.epilog:                                        ; preds = %entry, %sw.bb1, %sw.bb
+  ret void
+}
+
+declare spir_func i64 @_Z13get_global_idj(i32)
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpSwitch64.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpSwitch64.ll
new file mode 100644
index 0000000000000..5e4f1f14f8f91
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpSwitch64.ll
@@ -0,0 +1,63 @@
+;; __kernel void test_64(__global int* res)
+;; {
+;;     long tid = get_global_id(0);
+;;
+;;     switch(tid)
+;;     {
+;;     case 0:
+;;         res[tid] = 1;
+;;         break;
+;;     case 1:
+;;         res[tid] = 2;
+;;         break;
+;;     case 21474836481:
+;;         res[tid] = 3;
+;;         break;
+;;     }
+;; }
+;; bash$ clang -cc1 -triple spir64-unknown-unknown -x cl -cl-std=CL2.0 -O0 -include opencl.h -emit-llvm OpSwitch.cl -o test_64.ll
+
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: OpSwitch %[[#]] %[[#]] 0 0 %[[#]] 1 0 %[[#]] 1 5 %[[#]]
+
+define spir_kernel void @test_64(i32 addrspace(1)* %res) {
+entry:
+  %res.addr = alloca i32 addrspace(1)*, align 8
+  %tid = alloca i64, align 8
+  store i32 addrspace(1)* %res, i32 addrspace(1)** %res.addr, align 8
+  %call = call spir_func i64 @_Z13get_global_idj(i32 0)
+  store i64 %call, i64* %tid, align 8
+  %0 = load i64, i64* %tid, align 8
+  switch i64 %0, label %sw.epilog [
+    i64 0, label %sw.bb
+    i64 1, label %sw.bb1
+    i64 21474836481, label %sw.bb3
+  ]
+
+sw.bb:                                            ; preds = %entry
+  %1 = load i64, i64* %tid, align 8
+  %2 = load i32 addrspace(1)*, i32 addrspace(1)** %res.addr, align 8
+  %arrayidx = getelementptr inbounds i32, i32 addrspace(1)* %2, i64 %1
+  store i32 1, i32 addrspace(1)* %arrayidx, align 4
+  br label %sw.epilog
+
+sw.bb1:                                           ; preds = %entry
+  %3 = load i64, i64* %tid, align 8
+  %4 = load i32 addrspace(1)*, i32 addrspace(1)** %res.addr, align 8
+  %arrayidx2 = getelementptr inbounds i32, i32 addrspace(1)* %4, i64 %3
+  store i32 2, i32 addrspace(1)* %arrayidx2, align 4
+  br label %sw.epilog
+
+sw.bb3:                                           ; preds = %entry
+  %5 = load i64, i64* %tid, align 8
+  %6 = load i32 addrspace(1)*, i32 addrspace(1)** %res.addr, align 8
+  %arrayidx4 = getelementptr inbounds i32, i32 addrspace(1)* %6, i64 %5
+  store i32 3, i32 addrspace(1)* %arrayidx4, align 4
+  br label %sw.epilog
+
+sw.epilog:                                        ; preds = %entry, %sw.bb3, %sw.bb1, %sw.bb
+  ret void
+}
+
+declare spir_func i64 @_Z13get_global_idj(i32)
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpSwitchChar.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpSwitchChar.ll
new file mode 100644
index 0000000000000..2e1b01ef703fb
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpSwitchChar.ll
@@ -0,0 +1,51 @@
+;; __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: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: OpSwitch %[[#]] %[[#]] 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
+}
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/RelationalOperatorsFUnord.ll b/llvm/test/CodeGen/SPIRV/transcoding/RelationalOperatorsFUnord.ll
new file mode 100644
index 0000000000000..29164c20d95d2
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/RelationalOperatorsFUnord.ll
@@ -0,0 +1,70 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV:      %[[#bool:]] = OpTypeBool
+; CHECK-SPIRV:      %[[#bool2:]] = OpTypeVector %[[#bool]] 2
+
+; CHECK-SPIRV:      OpFunction
+; CHECK-SPIRV-NEXT: %[[#A:]] = OpFunctionParameter %[[#]]
+; CHECK-SPIRV-NEXT: %[[#B:]] = OpFunctionParameter %[[#]]
+; CHECK-SPIRV:      %[[#]] = OpFUnordEqual %[[#bool2]] %[[#A]] %[[#B]]
+; CHECK-SPIRV:      OpFunctionEnd
+
+@var = addrspace(1) global <2 x i1> zeroinitializer
+define spir_kernel void @testFUnordEqual(<2 x float> %a, <2 x float> %b) {
+entry:
+  %0 = fcmp ueq <2 x float> %a, %b
+  store <2 x i1> %0, <2 x i1> addrspace(1)* @var
+  ret void
+}
+
+; CHECK-SPIRV:      OpFunction
+; CHECK-SPIRV-NEXT: %[[#A:]] = OpFunctionParameter %[[#]]
+; CHECK-SPIRV-NEXT: %[[#B:]] = OpFunctionParameter %[[#]]
+; CHECK-SPIRV:      %[[#]] = OpFUnordGreaterThan %[[#bool2]] %[[#A]] %[[#B]]
+; CHECK-SPIRV:      OpFunctionEnd
+
+define spir_kernel void @testFUnordGreaterThan(<2 x float> %a, <2 x float> %b) {
+entry:
+  %0 = fcmp ugt <2 x float> %a, %b
+  store <2 x i1> %0, <2 x i1> addrspace(1)* @var
+  ret void
+}
+
+; CHECK-SPIRV:      OpFunction
+; CHECK-SPIRV-NEXT: %[[#A:]] = OpFunctionParameter %[[#]]
+; CHECK-SPIRV-NEXT: %[[#B:]] = OpFunctionParameter %[[#]]
+; CHECK-SPIRV:      %[[#]] = OpFUnordGreaterThanEqual %[[#bool2]] %[[#A]] %[[#B]]
+; CHECK-SPIRV:      OpFunctionEnd
+
+define spir_kernel void @testFUnordGreaterThanEqual(<2 x float> %a, <2 x float> %b) {
+entry:
+  %0 = fcmp uge <2 x float> %a, %b
+  store <2 x i1> %0, <2 x i1> addrspace(1)* @var
+  ret void
+}
+
+; CHECK-SPIRV:      OpFunction
+; CHECK-SPIRV-NEXT: %[[#A:]] = OpFunctionParameter %[[#]]
+; CHECK-SPIRV-NEXT: %[[#B:]] = OpFunctionParameter %[[#]]
+; CHECK-SPIRV:      %[[#]] = OpFUnordLessThan %[[#bool2]] %[[#A]] %[[#B]]
+; CHECK-SPIRV:      OpFunctionEnd
+
+define spir_kernel void @testFUnordLessThan(<2 x float> %a, <2 x float> %b) {
+entry:
+  %0 = fcmp ult <2 x float> %a, %b
+  store <2 x i1> %0, <2 x i1> addrspace(1)* @var
+  ret void
+}
+
+; CHECK-SPIRV:      OpFunction
+; CHECK-SPIRV-NEXT: %[[#A:]] = OpFunctionParameter %[[#]]
+; CHECK-SPIRV-NEXT: %[[#B:]] = OpFunctionParameter %[[#]]
+; CHECK-SPIRV:      %[[#]] = OpFUnordLessThanEqual %[[#bool2]] %[[#A]] %[[#B]]
+; CHECK-SPIRV:      OpFunctionEnd
+
+define spir_kernel void @testFUnordLessThanEqual(<2 x float> %a, <2 x float> %b) {
+entry:
+  %0 = fcmp ule <2 x float> %a, %b
+  store <2 x i1> %0, <2 x i1> addrspace(1)* @var
+  ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/TransFNeg.ll b/llvm/test/CodeGen/SPIRV/transcoding/TransFNeg.ll
new file mode 100644
index 0000000000000..eb52a775d38ee
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/TransFNeg.ll
@@ -0,0 +1,50 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: OpFNegate
+; CHECK-SPIRV: OpFNegate
+; CHECK-SPIRV: OpFNegate
+; CHECK-SPIRV: OpFNegate
+
+;; #pragma OPENCL EXTENSION cl_khr_fp64 : enable
+;; #pragma OPENCL EXTENSION cl_khr_fp16 : enable
+;;
+;; __kernel void foo(double a1, __global half *h, __global float *b0, __global double *b1, __global double8 *d) {
+;;    *h = -*h;
+;;    *b0 = -*b0;
+;;    *b1 = -a1;
+;;    *d = -*d;
+;; }
+
+define dso_local spir_kernel void @foo(double noundef %a1, half addrspace(1)* noundef %h, float addrspace(1)* noundef %b0, double addrspace(1)* noundef %b1, <8 x double> addrspace(1)* noundef %d) {
+entry:
+  %a1.addr = alloca double, align 8
+  %h.addr = alloca half addrspace(1)*, align 4
+  %b0.addr = alloca float addrspace(1)*, align 4
+  %b1.addr = alloca double addrspace(1)*, align 4
+  %d.addr = alloca <8 x double> addrspace(1)*, align 4
+  store double %a1, double* %a1.addr, align 8
+  store half addrspace(1)* %h, half addrspace(1)** %h.addr, align 4
+  store float addrspace(1)* %b0, float addrspace(1)** %b0.addr, align 4
+  store double addrspace(1)* %b1, double addrspace(1)** %b1.addr, align 4
+  store <8 x double> addrspace(1)* %d, <8 x double> addrspace(1)** %d.addr, align 4
+  %0 = load half addrspace(1)*, half addrspace(1)** %h.addr, align 4
+  %1 = load half, half addrspace(1)* %0, align 2
+  %fneg = fneg half %1
+  %2 = load half addrspace(1)*, half addrspace(1)** %h.addr, align 4
+  store half %fneg, half addrspace(1)* %2, align 2
+  %3 = load float addrspace(1)*, float addrspace(1)** %b0.addr, align 4
+  %4 = load float, float addrspace(1)* %3, align 4
+  %fneg1 = fneg float %4
+  %5 = load float addrspace(1)*, float addrspace(1)** %b0.addr, align 4
+  store float %fneg1, float addrspace(1)* %5, align 4
+  %6 = load double, double* %a1.addr, align 8
+  %fneg2 = fneg double %6
+  %7 = load double addrspace(1)*, double addrspace(1)** %b1.addr, align 4
+  store double %fneg2, double addrspace(1)* %7, align 8
+  %8 = load <8 x double> addrspace(1)*, <8 x double> addrspace(1)** %d.addr, align 4
+  %9 = load <8 x double>, <8 x double> addrspace(1)* %8, align 64
+  %fneg3 = fneg <8 x double> %9
+  %10 = load <8 x double> addrspace(1)*, <8 x double> addrspace(1)** %d.addr, align 4
+  store <8 x double> %fneg3, <8 x double> addrspace(1)* %10, align 64
+  ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/bitcast.ll b/llvm/test/CodeGen/SPIRV/transcoding/bitcast.ll
new file mode 100644
index 0000000000000..8dbf4d2c58b4b
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/bitcast.ll
@@ -0,0 +1,20 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+
+;; Check the bitcast is translated back to bitcast
+
+; CHECK: Bitcast
+
+define spir_kernel void @test_fn(<2 x i8> addrspace(1)* nocapture readonly %src, i16 addrspace(1)* nocapture %dst) {
+entry:
+  %call = tail call spir_func i64 @_Z13get_global_idj(i32 0)
+  %sext = shl i64 %call, 32
+  %idxprom = ashr exact i64 %sext, 32
+  %arrayidx = getelementptr inbounds <2 x i8>, <2 x i8> addrspace(1)* %src, i64 %idxprom
+  %0 = load <2 x i8>, <2 x i8> addrspace(1)* %arrayidx, align 2
+  %astype = bitcast <2 x i8> %0 to i16
+  %arrayidx2 = getelementptr inbounds i16, i16 addrspace(1)* %dst, i64 %idxprom
+  store i16 %astype, i16 addrspace(1)* %arrayidx2, align 2
+  ret void
+}
+
+declare spir_func i64 @_Z13get_global_idj(i32)
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/fneg.ll b/llvm/test/CodeGen/SPIRV/transcoding/fneg.ll
new file mode 100644
index 0000000000000..e17601a2c25a7
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/fneg.ll
@@ -0,0 +1,30 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: OpName %[[#r1:]] "r1"
+; CHECK-SPIRV: OpName %[[#r2:]] "r2"
+; CHECK-SPIRV: OpName %[[#r3:]] "r3"
+; CHECK-SPIRV: OpName %[[#r4:]] "r4"
+; CHECK-SPIRV: OpName %[[#r5:]] "r5"
+; CHECK-SPIRV: OpName %[[#r6:]] "r6"
+; CHECK-SPIRV: OpName %[[#r7:]] "r7"
+; CHECK-SPIRV-NOT: OpDecorate %{{.*}} FPFastMathMode
+; CHECK-SPIRV: %[[#float:]] = OpTypeFloat 32
+; CHECK-SPIRV: %[[#r1]] = OpFNegate %[[#float]]
+; CHECK-SPIRV: %[[#r2]] = OpFNegate %[[#float]]
+; CHECK-SPIRV: %[[#r3]] = OpFNegate %[[#float]]
+; CHECK-SPIRV: %[[#r4]] = OpFNegate %[[#float]]
+; CHECK-SPIRV: %[[#r5]] = OpFNegate %[[#float]]
+; CHECK-SPIRV: %[[#r6]] = OpFNegate %[[#float]]
+; CHECK-SPIRV: %[[#r7]] = OpFNegate %[[#float]]
+
+define spir_kernel void @testFNeg(float %a) local_unnamed_addr {
+entry:
+  %r1 = fneg float %a
+  %r2 = fneg nnan float %a
+  %r3 = fneg ninf float %a
+  %r4 = fneg nsz float %a
+  %r5 = fneg arcp float %a
+  %r6 = fneg fast float %a
+  %r7 = fneg nnan ninf float %a
+  ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/fp_contract_reassoc_fast_mode.ll b/llvm/test/CodeGen/SPIRV/transcoding/fp_contract_reassoc_fast_mode.ll
new file mode 100644
index 0000000000000..c035c35a339ee
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/fp_contract_reassoc_fast_mode.ll
@@ -0,0 +1,24 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV-NOT: OpCapability FPFastMathModeINTEL
+; CHECK-SPIRV:     OpName %[[#mu:]] "mul"
+; CHECK-SPIRV:     OpName %[[#su:]] "sub"
+; CHECK-SPIRV-NOT: OpDecorate %[[#mu]] FPFastMathMode AllowContractFastINTEL
+; CHECK-SPIRV-NOT: OpDecorate %[[#su]] FPFastMathMode AllowReassocINTEL
+
+define spir_kernel void @test(float %a, float %b) {
+entry:
+  %a.addr = alloca float, align 4
+  %b.addr = alloca float, align 4
+  store float %a, float* %a.addr, align 4
+  store float %b, float* %b.addr, align 4
+  %0 = load float, float* %a.addr, align 4
+  %1 = load float, float* %a.addr, align 4
+  %mul = fmul contract float %0, %1
+  store float %mul, float* %b.addr, align 4
+  %2 = load float, float* %b.addr, align 4
+  %3 = load float, float* %b.addr, align 4
+  %sub = fsub reassoc float %2, %3
+  store float %sub, float* %b.addr, align 4
+  ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/isequal.ll b/llvm/test/CodeGen/SPIRV/transcoding/isequal.ll
new file mode 100644
index 0000000000000..3c818afcdb167
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/isequal.ll
@@ -0,0 +1,22 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV-NOT: OpSConvert
+
+define spir_kernel void @math_kernel8(<8 x i32> addrspace(1)* nocapture %out, <8 x float> addrspace(1)* nocapture readonly %in1, <8 x float> addrspace(1)* nocapture readonly %in2) {
+entry:
+  %call = tail call spir_func i64 @_Z13get_global_idj(i32 0)
+  %sext = shl i64 %call, 32
+  %idxprom = ashr exact i64 %sext, 32
+  %arrayidx = getelementptr inbounds <8 x float>, <8 x float> addrspace(1)* %in1, i64 %idxprom
+  %0 = load <8 x float>, <8 x float> addrspace(1)* %arrayidx, align 32
+  %arrayidx2 = getelementptr inbounds <8 x float>, <8 x float> addrspace(1)* %in2, i64 %idxprom
+  %1 = load <8 x float>, <8 x float> addrspace(1)* %arrayidx2, align 32
+  %call3 = tail call spir_func <8 x i32> @_Z7isequalDv8_fDv8_f(<8 x float> %0, <8 x float> %1)
+  %arrayidx5 = getelementptr inbounds <8 x i32>, <8 x i32> addrspace(1)* %out, i64 %idxprom
+  store <8 x i32> %call3, <8 x i32> addrspace(1)* %arrayidx5, align 32
+  ret void
+}
+
+declare spir_func i64 @_Z13get_global_idj(i32)
+
+declare spir_func <8 x i32> @_Z7isequalDv8_fDv8_f(<8 x float>, <8 x float>)
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/memory_access.ll b/llvm/test/CodeGen/SPIRV/transcoding/memory_access.ll
new file mode 100644
index 0000000000000..fc757112f168c
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/memory_access.ll
@@ -0,0 +1,38 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV-NOT: OpStore %[[#]] %[[#]] Volatile Aligned 8
+; CHECK-SPIRV:     OpStore %[[#]] %[[#]] Volatile|Aligned 8
+; CHECK-SPIRV-NOT: %[[#]] = OpLoad %[[#]] %[[#]] Volatile Aligned 8
+; CHECK-SPIRV:     %[[#]] = OpLoad %[[#]] %[[#]] Volatile|Aligned 8
+; CHECK-SPIRV:     %[[#]] = OpLoad %[[#]] %[[#]] Aligned 4
+; CHECK-SPIRV-NOT: %[[#]] = OpLoad %[[#]] %[[#]] Volatile Aligned 8
+; CHECK-SPIRV:     %[[#]] = OpLoad %[[#]] %[[#]] Volatile|Aligned 8
+; CHECK-SPIRV-NOT: %[[#]] = OpLoad %[[#]] %[[#]] Volatile Aligned 0
+; CHECK-SPIRV:     %[[#]] = OpLoad %[[#]] %[[#]] Volatile|Aligned 8
+; CHECK-SPIRV-NOT: %[[#]] = OpLoad %[[#]] %[[#]] Volatile Aligned 8
+; CHECK-SPIRV:     %[[#]] = OpLoad %[[#]] %[[#]] Volatile|Aligned|Nontemporal 8
+; CHECK-SPIRV-NOT: OpStore %[[#]] %[[#]] Aligned 4
+; CHECK-SPIRV:     OpStore %[[#]] %[[#]] Aligned|Nontemporal 4
+; CHECK-SPIRV-NOT: OpStore %[[#]] %[[#]] Aligned 0
+; CHECK-SPIRV:     OpStore %[[#]] %[[#]]
+
+define spir_kernel void @test_load_store(i32 addrspace(1)* %destMemory, i32 addrspace(1)* %oldValues, i32 %newValue) {
+entry:
+  %ptr = alloca i32 addrspace(4)*, align 8
+  %0 = addrspacecast i32 addrspace(1)* %oldValues to i32 addrspace(4)*
+  store volatile i32 addrspace(4)* %0, i32 addrspace(4)** %ptr, align 8
+  %1 = load volatile i32 addrspace(4)*, i32 addrspace(4)** %ptr, align 8
+  %2 = load i32, i32 addrspace(4)* %1, align 4
+  %call = call spir_func i32 @_Z14atomic_cmpxchgPVU3AS1iii(i32 addrspace(1)* %destMemory, i32 %2, i32 %newValue)
+  %3 = load volatile i32 addrspace(4)*, i32 addrspace(4)** %ptr, align 8
+  %4 = load volatile i32 addrspace(4)*, i32 addrspace(4)** %ptr
+  %5 = load volatile i32 addrspace(4)*, i32 addrspace(4)** %ptr, align 8, !nontemporal !9
+  %arrayidx = getelementptr inbounds i32, i32 addrspace(4)* %3, i64 0
+  store i32 %call, i32 addrspace(4)* %arrayidx, align 4, !nontemporal !9
+  store i32 addrspace(4)* %5, i32 addrspace(4)** %ptr
+  ret void
+}
+
+declare spir_func i32 @_Z14atomic_cmpxchgPVU3AS1iii(i32 addrspace(1)*, i32, i32)
+
+!9 = !{i32 1}
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/non32.ll b/llvm/test/CodeGen/SPIRV/transcoding/non32.ll
new file mode 100644
index 0000000000000..d44e32141db75
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/non32.ll
@@ -0,0 +1,12 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+
+; CHECK: OpTypeInt 16
+; CHECK: OpIAdd
+
+define i16 @test_fn(i16 %arg0, i16 %arg1) {
+entry:
+  %0 = add i16 %arg0, %arg1
+  ret i16 %0
+}
+
+declare spir_func i64 @_Z13get_global_idj(i32)
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/vec_type_hint.ll b/llvm/test/CodeGen/SPIRV/transcoding/vec_type_hint.ll
new file mode 100644
index 0000000000000..1d512a894d540
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/vec_type_hint.ll
@@ -0,0 +1,51 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+;; kernel
+;; __attribute__((vec_type_hint(float4)))
+;; void test_float() {}
+
+;; kernel
+;; __attribute__((vec_type_hint(double)))
+;; void test_double() {}
+
+;; kernel
+;; __attribute__((vec_type_hint(uint4)))
+;; void test_uint() {}
+
+;; kernel
+;; __attribute__((vec_type_hint(int8)))
+;; void test_int() {}
+
+; CHECK-SPIRV: OpEntryPoint {{.*}} %[[#]] "test_float"
+; CHECK-SPIRV: OpEntryPoint {{.*}} %[[#]] "test_double"
+; CHECK-SPIRV: OpEntryPoint {{.*}} %[[#]] "test_uint"
+; CHECK-SPIRV: OpEntryPoint {{.*}} %[[#]] "test_int"
+; CHECK-SPIRV: OpExecutionMode %[[#]] VecTypeHint [[#]]
+; CHECK-SPIRV: OpExecutionMode %[[#]] VecTypeHint [[#]]
+; CHECK-SPIRV: OpExecutionMode %[[#]] VecTypeHint [[#]]
+; CHECK-SPIRV: OpExecutionMode %[[#]] VecTypeHint [[#]]
+
+define dso_local spir_kernel void @test_float() !vec_type_hint !4 {
+entry:
+  ret void
+}
+
+define dso_local spir_kernel void @test_double() !vec_type_hint !5 {
+entry:
+  ret void
+}
+
+define dso_local spir_kernel void @test_uint() !vec_type_hint !6 {
+entry:
+  ret void
+}
+
+define dso_local spir_kernel void @test_int() !vec_type_hint !7 {
+entry:
+  ret void
+}
+
+!4 = !{<4 x float> undef, i32 0}
+!5 = !{double undef, i32 0}
+!6 = !{<4 x i32> undef, i32 0}
+!7 = !{<8 x i32> undef, i32 1}
diff --git a/llvm/test/CodeGen/SPIRV/uitofp-with-bool.ll b/llvm/test/CodeGen/SPIRV/uitofp-with-bool.ll
new file mode 100644
index 0000000000000..75997a37bafff
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/uitofp-with-bool.ll
@@ -0,0 +1,132 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPV
+
+;; The IR was generated from the following source:
+;; void __kernel K(global float* A, int B) {
+;;   bool Cmp = B > 0;
+;;   A[0] = Cmp;
+;; }
+;; Command line:
+;; clang -x cl -cl-std=CL2.0 -target spir64 -emit-llvm -S -c test.cl
+
+
+; SPV-DAG: OpName %[[#s1:]] "s1"
+; SPV-DAG: OpName %[[#s2:]] "s2"
+; SPV-DAG: OpName %[[#s3:]] "s3"
+; SPV-DAG: OpName %[[#s4:]] "s4"
+; SPV-DAG: OpName %[[#s5:]] "s5"
+; SPV-DAG: OpName %[[#s6:]] "s6"
+; SPV-DAG: OpName %[[#s7:]] "s7"
+; SPV-DAG: OpName %[[#s8:]] "s8"
+; SPV-DAG: OpName %[[#z1:]] "z1"
+; SPV-DAG: OpName %[[#z2:]] "z2"
+; SPV-DAG: OpName %[[#z3:]] "z3"
+; SPV-DAG: OpName %[[#z4:]] "z4"
+; SPV-DAG: OpName %[[#z5:]] "z5"
+; SPV-DAG: OpName %[[#z6:]] "z6"
+; SPV-DAG: OpName %[[#z7:]] "z7"
+; SPV-DAG: OpName %[[#z8:]] "z8"
+; SPV-DAG: OpName %[[#ufp1:]] "ufp1"
+; SPV-DAG: OpName %[[#ufp2:]] "ufp2"
+; SPV-DAG: OpName %[[#sfp1:]] "sfp1"
+; SPV-DAG: OpName %[[#sfp2:]] "sfp2"
+; SPV-DAG: %[[#int_32:]] = OpTypeInt 32 0
+; SPV-DAG: %[[#int_8:]] = OpTypeInt 8 0
+; SPV-DAG: %[[#int_16:]] = OpTypeInt 16 0
+; SPV-DAG: %[[#int_64:]] = OpTypeInt 64 0
+; SPV-DAG: %[[#zero_32:]] = OpConstant %[[#int_32]] 0
+; SPV-DAG: %[[#one_32:]] = OpConstant %[[#int_32]] 1
+; SPV-DAG: %[[#zero_8:]] = OpConstantNull %[[#int_8]]
+; SPV-DAG: %[[#mone_8:]] = OpConstant %[[#int_8]] 255
+; SPV-DAG: %[[#zero_16:]] = OpConstantNull %[[#int_16]]
+; SPV-DAG: %[[#mone_16:]] = OpConstant %[[#int_16]] 65535
+; SPV-DAG: %[[#mone_32:]] = OpConstant %[[#int_32]] 4294967295
+; SPV-DAG: %[[#zero_64:]] = OpConstantNull %[[#int_64]]
+; SPV-DAG: %[[#mone_64:]] = OpConstant %[[#int_64]] 4294967295 4294967295
+; SPV-DAG: %[[#one_8:]] = OpConstant %[[#int_8]] 1
+; SPV-DAG: %[[#one_16:]] = OpConstant %[[#int_16]] 1
+; SPV-DAG: %[[#one_64:]] = OpConstant %[[#int_64]] 1 0
+; SPV-DAG: %[[#void:]] = OpTypeVoid
+; SPV-DAG: %[[#float:]] = OpTypeFloat 32
+; SPV-DAG: %[[#bool:]] = OpTypeBool
+; SPV-DAG: %[[#vec_8:]] = OpTypeVector %[[#int_8]] 2
+; SPV-DAG: %[[#vec_1:]] = OpTypeVector %[[#bool]] 2
+; SPV-DAG: %[[#vec_16:]] = OpTypeVector %[[#int_16]] 2
+; SPV-DAG: %[[#vec_32:]] = OpTypeVector %[[#int_32]] 2
+; SPV-DAG: %[[#vec_64:]] = OpTypeVector %[[#int_64]] 2
+; SPV-DAG: %[[#vec_float:]] = OpTypeVector %[[#float]] 2
+; SPV-DAG: %[[#zeros_8:]] = OpConstantNull %[[#vec_8]]
+; SPV-DAG: %[[#mones_8:]] = OpConstantComposite %[[#vec_8]] %[[#mone_8]] %[[#mone_8]]
+; SPV-DAG: %[[#zeros_16:]] = OpConstantNull %[[#vec_16]]
+; SPV-DAG: %[[#mones_16:]] = OpConstantComposite %[[#vec_16]] %[[#mone_16]] %[[#mone_16]]
+; SPV-DAG: %[[#zeros_32:]] = OpConstantNull %[[#vec_32]]
+; SPV-DAG: %[[#mones_32:]] = OpConstantComposite %[[#vec_32]] %[[#mone_32]] %[[#mone_32]]
+; SPV-DAG: %[[#zeros_64:]] = OpConstantNull %[[#vec_64]]
+; SPV-DAG: %[[#mones_64:]] = OpConstantComposite %[[#vec_64]] %[[#mone_64]] %[[#mone_64]]
+; SPV-DAG: %[[#ones_8:]] = OpConstantComposite %[[#vec_8]] %[[#one_8]] %[[#one_8]]
+; SPV-DAG: %[[#ones_16:]] = OpConstantComposite %[[#vec_16]] %[[#one_16]] %[[#one_16]]
+; SPV-DAG: %[[#ones_32:]] = OpConstantComposite %[[#vec_32]] %[[#one_32]] %[[#one_32]]
+; SPV-DAG: %[[#ones_64:]] = OpConstantComposite %[[#vec_64]] %[[#one_64]] %[[#one_64]]
+
+; SPV-DAG: OpFunction
+; SPV-DAG: %[[#A:]] = OpFunctionParameter %[[#]]
+; SPV-DAG: %[[#B:]] = OpFunctionParameter %[[#]]
+; SPV-DAG: %[[#i1s:]] = OpFunctionParameter %[[#]]
+; SPV-DAG: %[[#i1v:]] = OpFunctionParameter %[[#]]
+
+define dso_local spir_kernel void @K(float addrspace(1)* nocapture %A, i32 %B, i1 %i1s, <2 x i1> %i1v) local_unnamed_addr {
+entry:
+
+; SPV-DAG: %[[#cmp_res:]] = OpSGreaterThan %[[#bool]] %[[#B]] %[[#zero_32]]
+  %cmp = icmp sgt i32 %B, 0
+; SPV-DAG: %[[#select_res:]] = OpSelect %[[#int_32]] %[[#cmp_res]] %[[#one_32]] %[[#zero_32]]
+; SPV-DAG: %[[#utof_res:]] = OpConvertUToF %[[#float]] %[[#select_res]]
+  %conv = uitofp i1 %cmp to float
+; SPV-DAG: OpStore %[[#A]] %[[#utof_res]]
+  store float %conv, float addrspace(1)* %A, align 4;
+
+; SPV-DAG: %[[#s1]] = OpSelect %[[#int_8]] %[[#i1s]] %[[#mone_8]] %[[#zero_8]]
+  %s1 = sext i1 %i1s to i8
+; SPV-DAG: %[[#s2]] = OpSelect %[[#int_16]] %[[#i1s]] %[[#mone_16]] %[[#zero_16]]
+  %s2 = sext i1 %i1s to i16
+; SPV-DAG: %[[#s3]] = OpSelect %[[#int_32]] %[[#i1s]] %[[#mone_32]] %[[#zero_32]]
+  %s3 = sext i1 %i1s to i32
+; SPV-DAG: %[[#s4]] = OpSelect %[[#int_64]] %[[#i1s]] %[[#mone_64]] %[[#zero_64]]
+  %s4 = sext i1 %i1s to i64
+; SPV-DAG: %[[#s5]] = OpSelect %[[#vec_8]] %[[#i1v]] %[[#mones_8]] %[[#zeros_8]]
+  %s5 = sext <2 x i1> %i1v to <2 x i8>
+; SPV-DAG: %[[#s6]] = OpSelect %[[#vec_16]] %[[#i1v]] %[[#mones_16]] %[[#zeros_16]]
+  %s6 = sext <2 x i1> %i1v to <2 x i16>
+; SPV-DAG: %[[#s7]] = OpSelect %[[#vec_32]] %[[#i1v]] %[[#mones_32]] %[[#zeros_32]]
+  %s7 = sext <2 x i1> %i1v to <2 x i32>
+; SPV-DAG: %[[#s8]] = OpSelect %[[#vec_64]] %[[#i1v]] %[[#mones_64]] %[[#zeros_64]]
+  %s8 = sext <2 x i1> %i1v to <2 x i64>
+; SPV-DAG: %[[#z1]] = OpSelect %[[#int_8]] %[[#i1s]] %[[#one_8]] %[[#zero_8]]
+  %z1 = zext i1 %i1s to i8
+; SPV-DAG: %[[#z2]] = OpSelect %[[#int_16]] %[[#i1s]] %[[#one_16]] %[[#zero_16]]
+  %z2 = zext i1 %i1s to i16
+; SPV-DAG: %[[#z3]] = OpSelect %[[#int_32]] %[[#i1s]] %[[#one_32]] %[[#zero_32]]
+  %z3 = zext i1 %i1s to i32
+; SPV-DAG: %[[#z4]] = OpSelect %[[#int_64]] %[[#i1s]] %[[#one_64]] %[[#zero_64]]
+  %z4 = zext i1 %i1s to i64
+; SPV-DAG: %[[#z5]] = OpSelect %[[#vec_8]] %[[#i1v]] %[[#ones_8]] %[[#zeros_8]]
+  %z5 = zext <2 x i1> %i1v to <2 x i8>
+; SPV-DAG: %[[#z6]] = OpSelect %[[#vec_16]] %[[#i1v]] %[[#ones_16]] %[[#zeros_16]]
+  %z6 = zext <2 x i1> %i1v to <2 x i16>
+; SPV-DAG: %[[#z7]] = OpSelect %[[#vec_32]] %[[#i1v]] %[[#ones_32]] %[[#zeros_32]]
+  %z7 = zext <2 x i1> %i1v to <2 x i32>
+; SPV-DAG: %[[#z8]] = OpSelect %[[#vec_64]] %[[#i1v]] %[[#ones_64]] %[[#zeros_64]]
+  %z8 = zext <2 x i1> %i1v to <2 x i64>
+; SPV-DAG: %[[#ufp1_res:]] = OpSelect %[[#int_32]] %[[#i1s]] %[[#one_32]] %[[#zero_32]]
+; SPV-DAG: %[[#ufp1]] = OpConvertUToF %[[#float]] %[[#ufp1_res]]
+  %ufp1 = uitofp i1 %i1s to float
+; SPV-DAG: %[[#ufp2_res:]] = OpSelect %[[#vec_32]] %[[#i1v]] %[[#ones_32]] %[[#zeros_32]]
+; SPV-DAG: %[[#ufp2]] = OpConvertUToF %[[#vec_float]] %[[#ufp2_res]]
+  %ufp2 = uitofp <2 x i1> %i1v to <2 x float>
+; SPV-DAG: %[[#sfp1_res:]] = OpSelect %[[#int_32]] %[[#i1s]] %[[#one_32]] %[[#zero_32]]
+; SPV-DAG: %[[#sfp1]] = OpConvertSToF %[[#float]] %[[#sfp1_res]]
+  %sfp1 = sitofp i1 %i1s to float
+; SPV-DAG: %[[#sfp2_res:]] = OpSelect %[[#vec_32]] %[[#i1v]] %[[#ones_32]] %[[#zeros_32]]
+; SPV-DAG: %[[#sfp2]] = OpConvertSToF %[[#vec_float]] %[[#sfp2_res]]
+  %sfp2 = sitofp <2 x i1> %i1v to <2 x float>
+  ret void
+}