[llvm] 8468e67 - [SPIRV] Add tests to improve test coverage
Andrey Tretyakov via llvm-commits
llvm-commits at lists.llvm.org
Tue Aug 2 10:24:14 PDT 2022
Author: Andrey Tretyakov
Date: 2022-08-02T20:22:40+03:00
New Revision: 8468e67495eda83e2490224be46967ddc29821d2
URL: https://github.com/llvm/llvm-project/commit/8468e67495eda83e2490224be46967ddc29821d2
DIFF: https://github.com/llvm/llvm-project/commit/8468e67495eda83e2490224be46967ddc29821d2.diff
LOG: [SPIRV] Add tests to improve test coverage
Differential Revision: https://reviews.llvm.org/D130597
Added:
llvm/test/CodeGen/SPIRV/LinkOnceODR.ll
llvm/test/CodeGen/SPIRV/OpVectorInsertDynamic.ll
llvm/test/CodeGen/SPIRV/atomicrmw.ll
llvm/test/CodeGen/SPIRV/exec_mode_float_control_khr.ll
llvm/test/CodeGen/SPIRV/instructions/call-trivial-function.ll
llvm/test/CodeGen/SPIRV/instructions/scalar-bitwise-operations.ll
llvm/test/CodeGen/SPIRV/instructions/vector-bitwise-operations.ll
llvm/test/CodeGen/SPIRV/instructions/vector-shuffle.ll
llvm/test/CodeGen/SPIRV/linked-list.ll
llvm/test/CodeGen/SPIRV/llvm-intrinsics/assume.ll
llvm/test/CodeGen/SPIRV/llvm-intrinsics/expect.ll
llvm/test/CodeGen/SPIRV/llvm-intrinsics/fshr.ll
llvm/test/CodeGen/SPIRV/llvm-intrinsics/invariant.ll
llvm/test/CodeGen/SPIRV/llvm-intrinsics/umul.with.overflow.ll
llvm/test/CodeGen/SPIRV/memory_model_md.ll
llvm/test/CodeGen/SPIRV/multi_md.ll
llvm/test/CodeGen/SPIRV/opencl/basic/get_global_offset.ll
llvm/test/CodeGen/SPIRV/opencl/basic/progvar_prog_scope_init.ll
llvm/test/CodeGen/SPIRV/opencl/basic/progvar_prog_scope_uninit.ll
llvm/test/CodeGen/SPIRV/preprocess-metadata.ll
llvm/test/CodeGen/SPIRV/pstruct.ll
llvm/test/CodeGen/SPIRV/sitofp-with-bool.ll
llvm/test/CodeGen/SPIRV/spirv_param_decorations.ll
llvm/test/CodeGen/SPIRV/spirv_param_decorations_quals.ll
llvm/test/CodeGen/SPIRV/store.ll
llvm/test/CodeGen/SPIRV/struct.ll
llvm/test/CodeGen/SPIRV/transcoding/ConvertPtr.ll
llvm/test/CodeGen/SPIRV/transcoding/DecorationAlignment.ll
llvm/test/CodeGen/SPIRV/transcoding/DivRem.ll
llvm/test/CodeGen/SPIRV/transcoding/ExecutionMode_SPIR_to_SPIRV.ll
llvm/test/CodeGen/SPIRV/transcoding/GlobalFunAnnotate.ll
llvm/test/CodeGen/SPIRV/transcoding/OpConstantBool.ll
llvm/test/CodeGen/SPIRV/transcoding/OpPhi_ArgumentsPlaceholders.ll
llvm/test/CodeGen/SPIRV/transcoding/OpSwitch32.ll
llvm/test/CodeGen/SPIRV/transcoding/OpSwitch64.ll
llvm/test/CodeGen/SPIRV/transcoding/OpSwitchChar.ll
llvm/test/CodeGen/SPIRV/transcoding/RelationalOperatorsFUnord.ll
llvm/test/CodeGen/SPIRV/transcoding/TransFNeg.ll
llvm/test/CodeGen/SPIRV/transcoding/bitcast.ll
llvm/test/CodeGen/SPIRV/transcoding/fneg.ll
llvm/test/CodeGen/SPIRV/transcoding/fp_contract_reassoc_fast_mode.ll
llvm/test/CodeGen/SPIRV/transcoding/isequal.ll
llvm/test/CodeGen/SPIRV/transcoding/memory_access.ll
llvm/test/CodeGen/SPIRV/transcoding/non32.ll
llvm/test/CodeGen/SPIRV/transcoding/vec_type_hint.ll
llvm/test/CodeGen/SPIRV/uitofp-with-bool.ll
Modified:
Removed:
################################################################################
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
+
+ at 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
+
+ at ui = common dso_local addrspace(1) global i32 0, align 4
+ at 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
+
+ at 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 }
+
+ at 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]]
+ at var = addrspace(1) global i8 0, align 1
+ at g_var = addrspace(1) global i8 1, align 1
+ at a_var = addrspace(1) global [2 x i8] c"\01\01", align 1
+ at 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]]
+ at var = addrspace(1) global <2 x i8> zeroinitializer, align 2
+ at g_var = addrspace(1) global <2 x i8> zeroinitializer, align 2
+ at a_var = addrspace(1) global [2 x <2 x i8>] zeroinitializer, align 2
+ at 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"
+
+ at .str = private unnamed_addr constant [23 x i8] c"annotation_on_function\00", section "llvm.metadata"
+ at .str.1 = private unnamed_addr constant [6 x i8] c"an.cl\00", section "llvm.metadata"
+ at 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
+
+ at 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
+}
More information about the llvm-commits
mailing list