[llvm] f20c9c4 - [SPIRV] Add tests to improve test coverage

[Andrey Tretyakov via llvm-commits]

Author: Andrey Tretyakov
Date: 2022-09-02T13:19:28+03:00
New Revision: f20c9c42d2ca13354c8aebc093a61ce692903eed

URL: https://github.com/llvm/llvm-project/commit/f20c9c42d2ca13354c8aebc093a61ce692903eed
DIFF: https://github.com/llvm/llvm-project/commit/f20c9c42d2ca13354c8aebc093a61ce692903eed.diff

LOG: [SPIRV] Add tests to improve test coverage

Differential Revision: https://reviews.llvm.org/D132903

Added: 
    llvm/test/CodeGen/SPIRV/SampledImageRetType.ll
    llvm/test/CodeGen/SPIRV/event_no_group_cap.ll
    llvm/test/CodeGen/SPIRV/image-unoptimized.ll
    llvm/test/CodeGen/SPIRV/image.ll
    llvm/test/CodeGen/SPIRV/image_decl_func_arg.ll
    llvm/test/CodeGen/SPIRV/image_dim.ll
    llvm/test/CodeGen/SPIRV/image_store.ll
    llvm/test/CodeGen/SPIRV/instructions/float-fast-flags.ll
    llvm/test/CodeGen/SPIRV/instructions/scalar-floating-point-arithmetic.ll
    llvm/test/CodeGen/SPIRV/link-attribute.ll
    llvm/test/CodeGen/SPIRV/llvm-intrinsics/abs.ll
    llvm/test/CodeGen/SPIRV/llvm-intrinsics/ceil.ll
    llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctlz.ll
    llvm/test/CodeGen/SPIRV/llvm-intrinsics/cttz.ll
    llvm/test/CodeGen/SPIRV/llvm-intrinsics/fabs.ll
    llvm/test/CodeGen/SPIRV/llvm-intrinsics/nearbyint.ll
    llvm/test/CodeGen/SPIRV/llvm-intrinsics/sqrt.ll
    llvm/test/CodeGen/SPIRV/mangled_function.ll
    llvm/test/CodeGen/SPIRV/no_capability_shader.ll
    llvm/test/CodeGen/SPIRV/opencl.queue_t.ll
    llvm/test/CodeGen/SPIRV/read_image.ll
    llvm/test/CodeGen/SPIRV/spirv.Queue.ll
    llvm/test/CodeGen/SPIRV/transcoding/OpConstantSampler.ll
    llvm/test/CodeGen/SPIRV/transcoding/OpGroupAsyncCopy.ll
    llvm/test/CodeGen/SPIRV/transcoding/OpImageQuerySize.ll
    llvm/test/CodeGen/SPIRV/transcoding/OpImageReadMS.ll
    llvm/test/CodeGen/SPIRV/transcoding/OpImageSampleExplicitLod.ll
    llvm/test/CodeGen/SPIRV/transcoding/OpImageWrite.ll
    llvm/test/CodeGen/SPIRV/transcoding/OpSwitchUnreachable.ll
    llvm/test/CodeGen/SPIRV/transcoding/SampledImage.ll
    llvm/test/CodeGen/SPIRV/transcoding/check_ro_qualifier.ll
    llvm/test/CodeGen/SPIRV/transcoding/cl-types.ll
    llvm/test/CodeGen/SPIRV/transcoding/clk_event_t.ll
    llvm/test/CodeGen/SPIRV/transcoding/explicit-conversions.ll
    llvm/test/CodeGen/SPIRV/transcoding/image_get_size_with_access_qualifiers.ll
    llvm/test/CodeGen/SPIRV/transcoding/image_with_access_qualifiers.ll
    llvm/test/CodeGen/SPIRV/transcoding/optional-core-features-multiple.ll
    llvm/test/CodeGen/SPIRV/transcoding/spirv-types.ll

Modified: 
    

Removed: 
    


################################################################################
diff  --git a/llvm/test/CodeGen/SPIRV/SampledImageRetType.ll b/llvm/test/CodeGen/SPIRV/SampledImageRetType.ll
new file mode 100644
index 0000000000000..3bdda3655eaad
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/SampledImageRetType.ll
@@ -0,0 +1,29 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+
+%opencl.image1d_ro_t = type opaque
+; CHECK: %[[#image1d_t:]] = OpTypeImage
+%opencl.sampler_t = type opaque
+; CHECK: %[[#sampler_t:]] = OpTypeSampler
+; CHECK: %[[#sampled_image_t:]] = OpTypeSampledImage
+
+declare dso_local spir_func i8 addrspace(4)* @_Z20__spirv_SampledImageI14ocl_image1d_roPvET0_T_11ocl_sampler(%opencl.image1d_ro_t addrspace(1)*, %opencl.sampler_t addrspace(2)*) local_unnamed_addr
+
+declare dso_local spir_func <4 x float> @_Z30__spirv_ImageSampleExplicitLodIPvDv4_fiET0_T_T1_if(i8 addrspace(4)*, i32, i32, float) local_unnamed_addr
+
+ at __spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(2) constant <3 x i64>, align 32
+
+define weak_odr dso_local spir_kernel void @_ZTS17image_kernel_readILi1EE(%opencl.image1d_ro_t addrspace(1)*, %opencl.sampler_t addrspace(2)*) {
+; CHECK: OpFunction
+; CHECK: %[[#image:]] = OpFunctionParameter %[[#image1d_t]]
+; CHECK: %[[#sampler:]] = OpFunctionParameter %[[#sampler_t]]
+  %3 = load <3 x i64>, <3 x i64> addrspace(2)* @__spirv_BuiltInGlobalInvocationId, align 32
+  %4 = extractelement <3 x i64> %3, i64 0
+  %5 = trunc i64 %4 to i32
+  %6 = tail call spir_func i8 addrspace(4)* @_Z20__spirv_SampledImageI14ocl_image1d_roPvET0_T_11ocl_sampler(%opencl.image1d_ro_t addrspace(1)* %0, %opencl.sampler_t addrspace(2)* %1)
+  %7 = tail call spir_func <4 x float> @_Z30__spirv_ImageSampleExplicitLodIPvDv4_fiET0_T_T1_if(i8 addrspace(4)* %6, i32 %5, i32 2, float 0.000000e+00)
+
+; CHECK: %[[#sampled_image:]] = OpSampledImage %[[#sampled_image_t]] %[[#image]] %[[#sampler]]
+; CHECK: %[[#]] = OpImageSampleExplicitLod %[[#]] %[[#sampled_image]] %[[#]] {{.*}} %[[#]]
+
+  ret void
+}

diff  --git a/llvm/test/CodeGen/SPIRV/event_no_group_cap.ll b/llvm/test/CodeGen/SPIRV/event_no_group_cap.ll
new file mode 100644
index 0000000000000..af4ac7cf44500
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/event_no_group_cap.ll
@@ -0,0 +1,21 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+
+; __kernel void test_fn( const __global char *src)
+; {
+;     wait_group_events(0, NULL);
+; }
+
+; CHECK-NOT: OpCapability Groups
+; CHECK: OpGroupWaitEvents
+
+%opencl.event_t = type opaque
+
+define dso_local spir_kernel void @test_fn(i8 addrspace(1)* noundef %src) {
+entry:
+  %src.addr = alloca i8 addrspace(1)*, align 8
+  store i8 addrspace(1)* %src, i8 addrspace(1)** %src.addr, align 8
+  call spir_func void @_Z17wait_group_eventsiPU3AS49ocl_event(i32 noundef 0, %opencl.event_t* addrspace(4)* noundef null)
+  ret void
+}
+
+declare spir_func void @_Z17wait_group_eventsiPU3AS49ocl_event(i32 noundef, %opencl.event_t* addrspace(4)* noundef)

diff  --git a/llvm/test/CodeGen/SPIRV/image-unoptimized.ll b/llvm/test/CodeGen/SPIRV/image-unoptimized.ll
new file mode 100644
index 0000000000000..3604f1bddb188
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/image-unoptimized.ll
@@ -0,0 +1,89 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
+
+; CHECK:     %[[#TypeImage:]] = OpTypeImage
+; CHECK:     %[[#TypeSampler:]] = OpTypeSampler
+; CHECK-DAG: %[[#TypeImagePtr:]] = OpTypePointer {{.*}} %[[#TypeImage]]
+; CHECK-DAG: %[[#TypeSamplerPtr:]] = OpTypePointer {{.*}} %[[#TypeSampler]]
+
+; CHECK:     %[[#srcimg:]] = OpFunctionParameter %[[#TypeImage]]
+; CHECK:     %[[#sampler:]] = OpFunctionParameter %[[#TypeSampler]]
+
+; CHECK:     %[[#srcimg_addr:]] = OpVariable %[[#TypeImagePtr]]
+; CHECK:     %[[#sampler_addr:]] = OpVariable %[[#TypeSamplerPtr]]
+
+; CHECK:     OpStore %[[#srcimg_addr]] %[[#srcimg]]
+; CHECK:     OpStore %[[#sampler_addr]] %[[#sampler]]
+
+; CHECK:     %[[#srcimg_val:]] = OpLoad %[[#]] %[[#srcimg_addr]]
+; CHECK:     %[[#sampler_val:]] = OpLoad %[[#]] %[[#sampler_addr]]
+
+; CHECK:     %[[#]] = OpSampledImage %[[#]] %[[#srcimg_val]] %[[#sampler_val]]
+; CHECK-NEXT: OpImageSampleExplicitLod
+
+; CHECK:     %[[#srcimg_val:]] = OpLoad %[[#]] %[[#srcimg_addr]]
+; CHECK:     %[[#]] = OpImageQuerySizeLod %[[#]] %[[#srcimg_val]]
+
+;; Excerpt from opencl-c-base.h
+;; typedef float float4 __attribute__((ext_vector_type(4)));
+;; typedef int int2 __attribute__((ext_vector_type(2)));
+;; typedef __SIZE_TYPE__ size_t;
+;;
+;; Excerpt from opencl-c.h to speed up compilation.
+;; #define __ovld __attribute__((overloadable))
+;; #define __purefn __attribute__((pure))
+;; #define __cnfn __attribute__((const))
+;; size_t __ovld __cnfn get_global_id(unsigned int dimindx);
+;; int __ovld __cnfn get_image_width(read_only image2d_t image);
+;; float4 __purefn __ovld read_imagef(read_only image2d_t image, sampler_t sampler, int2 coord);
+;;
+;;
+;; __kernel void test_fn(image2d_t srcimg, sampler_t sampler, global float4 *results) {
+;;   int tid_x = get_global_id(0);
+;;   int tid_y = get_global_id(1);
+;;   results[tid_x + tid_y * get_image_width(srcimg)] = read_imagef(srcimg, sampler, (int2){tid_x, tid_y});
+;; }
+
+%opencl.image2d_ro_t = type opaque
+%opencl.sampler_t = type opaque
+
+define dso_local spir_kernel void @test_fn(%opencl.image2d_ro_t addrspace(1)* %srcimg, %opencl.sampler_t addrspace(2)* %sampler, <4 x float> addrspace(1)* noundef %results) {
+entry:
+  %srcimg.addr = alloca %opencl.image2d_ro_t addrspace(1)*, align 4
+  %sampler.addr = alloca %opencl.sampler_t addrspace(2)*, align 4
+  %results.addr = alloca <4 x float> addrspace(1)*, align 4
+  %tid_x = alloca i32, align 4
+  %tid_y = alloca i32, align 4
+  %.compoundliteral = alloca <2 x i32>, align 8
+  store %opencl.image2d_ro_t addrspace(1)* %srcimg, %opencl.image2d_ro_t addrspace(1)** %srcimg.addr, align 4
+  store %opencl.sampler_t addrspace(2)* %sampler, %opencl.sampler_t addrspace(2)** %sampler.addr, align 4
+  store <4 x float> addrspace(1)* %results, <4 x float> addrspace(1)** %results.addr, align 4
+  %call = call spir_func i32 @_Z13get_global_idj(i32 noundef 0)
+  store i32 %call, i32* %tid_x, align 4
+  %call1 = call spir_func i32 @_Z13get_global_idj(i32 noundef 1)
+  store i32 %call1, i32* %tid_y, align 4
+  %0 = load %opencl.image2d_ro_t addrspace(1)*, %opencl.image2d_ro_t addrspace(1)** %srcimg.addr, align 4
+  %1 = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** %sampler.addr, align 4
+  %2 = load i32, i32* %tid_x, align 4
+  %vecinit = insertelement <2 x i32> undef, i32 %2, i32 0
+  %3 = load i32, i32* %tid_y, align 4
+  %vecinit2 = insertelement <2 x i32> %vecinit, i32 %3, i32 1
+  store <2 x i32> %vecinit2, <2 x i32>* %.compoundliteral, align 8
+  %4 = load <2 x i32>, <2 x i32>* %.compoundliteral, align 8
+  %call3 = call spir_func <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv2_i(%opencl.image2d_ro_t addrspace(1)* %0, %opencl.sampler_t addrspace(2)* %1, <2 x i32> noundef %4)
+  %5 = load <4 x float> addrspace(1)*, <4 x float> addrspace(1)** %results.addr, align 4
+  %6 = load i32, i32* %tid_x, align 4
+  %7 = load i32, i32* %tid_y, align 4
+  %8 = load %opencl.image2d_ro_t addrspace(1)*, %opencl.image2d_ro_t addrspace(1)** %srcimg.addr, align 4
+  %call4 = call spir_func i32 @_Z15get_image_width14ocl_image2d_ro(%opencl.image2d_ro_t addrspace(1)* %8)
+  %mul = mul nsw i32 %7, %call4
+  %add = add nsw i32 %6, %mul
+  %arrayidx = getelementptr inbounds <4 x float>, <4 x float> addrspace(1)* %5, i32 %add
+  store <4 x float> %call3, <4 x float> addrspace(1)* %arrayidx, align 16
+  ret void
+}
+
+declare spir_func i32 @_Z13get_global_idj(i32 noundef)
+
+declare spir_func <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv2_i(%opencl.image2d_ro_t addrspace(1)*, %opencl.sampler_t addrspace(2)*, <2 x i32> noundef)
+
+declare spir_func i32 @_Z15get_image_width14ocl_image2d_ro(%opencl.image2d_ro_t addrspace(1)*)

diff  --git a/llvm/test/CodeGen/SPIRV/image.ll b/llvm/test/CodeGen/SPIRV/image.ll
new file mode 100644
index 0000000000000..dec605d152ca2
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/image.ll
@@ -0,0 +1,31 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV:     %[[#VOID_TY:]] = OpTypeVoid
+; CHECK-SPIRV-DAG: %[[#]] = OpTypeImage %[[#VOID_TY]] 2D 0 0 0 0 Unknown ReadOnly
+; CHECK-SPIRV-DAG: %[[#]] = OpTypeImage %[[#VOID_TY]] 2D 0 0 0 0 Unknown WriteOnly
+; CHECK-SPIRV-NOT: %[[#]] = OpTypeImage %[[#VOID_TY]] 2D 0 0 0 0 Unknown ReadOnly
+; CHECK-SPIRV:     OpImageSampleExplicitLod
+; CHECK-SPIRV:     OpImageWrite
+
+%opencl.image2d_t = type opaque
+
+define spir_kernel void @image_copy(%opencl.image2d_t addrspace(1)* readnone %image1, %opencl.image2d_t addrspace(1)* %image2) !kernel_arg_access_qual !1 {
+entry:
+  %call = tail call spir_func i64 @_Z13get_global_idj(i32 0)
+  %conv = trunc i64 %call to i32
+  %call1 = tail call spir_func i64 @_Z13get_global_idj(i32 1)
+  %conv2 = trunc i64 %call1 to i32
+  %vecinit = insertelement <2 x i32> undef, i32 %conv, i32 0
+  %vecinit3 = insertelement <2 x i32> %vecinit, i32 %conv2, i32 1
+  %call4 = tail call spir_func <4 x float> @_Z11read_imagef11ocl_image2d11ocl_samplerDv2_i(%opencl.image2d_t addrspace(1)* %image1, i32 20, <2 x i32> %vecinit3)
+  tail call spir_func void @_Z12write_imagef11ocl_image2dDv2_iDv4_f(%opencl.image2d_t addrspace(1)* %image2, <2 x i32> %vecinit3, <4 x float> %call4)
+  ret void
+}
+
+declare spir_func i64 @_Z13get_global_idj(i32)
+
+declare spir_func <4 x float> @_Z11read_imagef11ocl_image2d11ocl_samplerDv2_i(%opencl.image2d_t addrspace(1)*, i32, <2 x i32>)
+
+declare spir_func void @_Z12write_imagef11ocl_image2dDv2_iDv4_f(%opencl.image2d_t addrspace(1)*, <2 x i32>, <4 x float>)
+
+!1 = !{!"read_only", !"write_only"}

diff  --git a/llvm/test/CodeGen/SPIRV/image_decl_func_arg.ll b/llvm/test/CodeGen/SPIRV/image_decl_func_arg.ll
new file mode 100644
index 0000000000000..ae8ded7713421
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/image_decl_func_arg.ll
@@ -0,0 +1,28 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV:     %[[#TypeImage:]] = OpTypeImage
+; CHECK-SPIRV-NOT: OpTypeImage
+; CHECK-SPIRV:     %[[#]] = OpTypeFunction %[[#]] %[[#TypeImage]]
+; CHECK-SPIRV:     %[[#]] = OpTypeFunction %[[#]] %[[#TypeImage]]
+; CHECK-SPIRV:     %[[#]] = OpFunctionParameter %[[#TypeImage]]
+; CHECK-SPIRV:     %[[#]] = OpFunctionParameter %[[#TypeImage]]
+; CHECK-SPIRV:     %[[#ParamID:]] = OpFunctionParameter %[[#TypeImage]]
+; CHECK-SPIRV:     %[[#]] = OpFunctionCall %[[#]] %[[#]] %[[#ParamID]]
+
+%opencl.image2d_ro_t = type opaque
+
+define spir_func void @f0(%opencl.image2d_ro_t addrspace(1)* %v2, <2 x float> %v3) {
+entry:
+  ret void
+}
+
+define spir_func void @f1(%opencl.image2d_ro_t addrspace(1)* %v2, <2 x float> %v3) {
+entry:
+  ret void
+}
+
+define spir_kernel void @test(%opencl.image2d_ro_t addrspace(1)* %v1) {
+entry:
+  call spir_func void @f0(%opencl.image2d_ro_t addrspace(1)* %v1, <2 x float> <float 1.000000e+00, float 5.000000e+00>)
+  ret void
+}

diff  --git a/llvm/test/CodeGen/SPIRV/image_dim.ll b/llvm/test/CodeGen/SPIRV/image_dim.ll
new file mode 100644
index 0000000000000..45353718d6a6e
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/image_dim.ll
@@ -0,0 +1,12 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV-DAG: OpCapability Sampled1D
+; CHECK-SPIRV-DAG: OpCapability SampledBuffer
+
+%opencl.image1d_t = type opaque
+%opencl.image1d_buffer_t = type opaque
+
+define spir_kernel void @image_d(%opencl.image1d_t addrspace(1)* %image1d_td6, %opencl.image1d_buffer_t addrspace(1)* %image1d_buffer_td8) {
+entry:
+  ret void
+}

diff  --git a/llvm/test/CodeGen/SPIRV/image_store.ll b/llvm/test/CodeGen/SPIRV/image_store.ll
new file mode 100644
index 0000000000000..d87853041b62c
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/image_store.ll
@@ -0,0 +1,23 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+
+;; Image types may be represented in two ways while translating to SPIR-V:
+;; - OpenCL form, for example, '%opencl.image2d_ro_t',
+;; - SPIR-V form, for example, '%spirv.Image._void_1_0_0_0_0_0_0',
+;; but it is still one type which should be translated to one SPIR-V type.
+;;
+;; The test checks that the code below is successfully translated and only one
+;; SPIR-V type for images is generated.
+
+; CHECK:     OpTypeImage
+; CHECK-NOT: OpTypeImage
+
+%opencl.image2d_ro_t = type opaque
+%spirv.Image._void_1_0_0_0_0_0_0 = type opaque
+
+define spir_kernel void @read_image(%opencl.image2d_ro_t addrspace(1)* %srcimg) {
+entry:
+  %srcimg.addr = alloca %opencl.image2d_ro_t addrspace(1)*, align 8
+  %spirvimg.addr = alloca %spirv.Image._void_1_0_0_0_0_0_0 addrspace(1)*, align 8
+  store %opencl.image2d_ro_t addrspace(1)* %srcimg, %opencl.image2d_ro_t addrspace(1)** %srcimg.addr, align 8
+  ret void
+}

diff  --git a/llvm/test/CodeGen/SPIRV/instructions/float-fast-flags.ll b/llvm/test/CodeGen/SPIRV/instructions/float-fast-flags.ll
new file mode 100644
index 0000000000000..1db5f8bb0ee36
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/instructions/float-fast-flags.ll
@@ -0,0 +1,95 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
+
+; DISABLED-CHECK-DAG: OpName [[FNEG:%.+]] "scalar_fneg"
+; CHECK-DAG: OpName [[FADD:%.+]] "test_fadd"
+; CHECK-DAG: OpName [[FSUB:%.+]] "test_fsub"
+; CHECK-DAG: OpName [[FMUL:%.+]] "test_fmul"
+; CHECK-DAG: OpName [[FDIV:%.+]] "test_fdiv"
+; CHECK-DAG: OpName [[FREM:%.+]] "test_frem"
+; CHECK-DAG: OpName [[FMA:%.+]] "test_fma"
+
+; CHECK-DAG: [[F32Ty:%.+]] = OpTypeFloat 32
+; CHECK-DAG: [[FNTy:%.+]] = OpTypeFunction [[F32Ty]] [[F32Ty]] [[F32Ty]]
+
+
+; CHECK:      [[FADD]] = OpFunction [[F32Ty]] None [[FNTy]]
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[F32Ty]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[F32Ty]]
+; CHECK-NEXT: OpLabel
+; CHECK-NEXT: [[C:%.+]] = OpFAdd [[F32Ty]] [[A]] [[B]]
+;; TODO: OpDecorate checks
+; CHECK-NEXT: OpReturnValue [[C]]
+; CHECK-NEXT: OpFunctionEnd
+define float @test_fadd(float %a, float %b) {
+    %c = fadd nnan ninf float %a, %b
+    ret float %c
+}
+
+; CHECK:      [[FSUB]] = OpFunction [[F32Ty]] None [[FNTy]]
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[F32Ty]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[F32Ty]]
+; CHECK-NEXT: OpLabel
+; CHECK-NEXT: [[C:%.+]] = OpFSub [[F32Ty]] [[A]] [[B]]
+;; TODO: OpDecorate checks
+; CHECK-NEXT: OpReturnValue [[C]]
+; CHECK-NEXT: OpFunctionEnd
+define float @test_fsub(float %a, float %b) {
+    %c = fsub fast float %a, %b
+    ret float %c
+}
+
+; CHECK:      [[FMUL]] = OpFunction [[F32Ty]] None [[FNTy]]
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[F32Ty]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[F32Ty]]
+; CHECK-NEXT: OpLabel
+; CHECK-NEXT: [[C:%.+]] = OpFMul [[F32Ty]] [[A]] [[B]]
+;; TODO: OpDecorate checks]
+; CHECK-NEXT: OpReturnValue [[C]]
+; CHECK-NEXT: OpFunctionEnd
+define float @test_fmul(float %a, float %b) {
+    %c = fmul contract float %a, %b
+    ret float %c
+}
+
+; CHECK:      [[FDIV]] = OpFunction [[F32Ty]] None [[FNTy]]
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[F32Ty]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[F32Ty]]
+; CHECK-NEXT: OpLabel
+; CHECK-NEXT: [[C:%.+]] = OpFDiv [[F32Ty]] [[A]] [[B]]
+;; TODO: OpDecorate checks
+; CHECK-NEXT: OpReturnValue [[C]]
+; CHECK-NEXT: OpFunctionEnd
+define float @test_fdiv(float %a, float %b) {
+    %c = fdiv arcp nsz float %a, %b
+    ret float %c
+}
+
+; CHECK:      [[FREM]] = OpFunction [[F32Ty]] None [[FNTy]]
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[F32Ty]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[F32Ty]]
+; CHECK-NEXT: OpLabel
+; CHECK-NEXT: [[C:%.+]] = OpFRem [[F32Ty]] [[A]] [[B]]
+;; TODO: OpDecorate checks
+; CHECK-NEXT: OpReturnValue [[C]]
+; CHECK-NEXT: OpFunctionEnd
+define float @test_frem(float %a, float %b) {
+    %c = frem nsz float %a, %b
+    ret float %c
+}
+
+
+declare float @llvm.fma.f32(float, float, float)
+
+; CHECK:      [[FMA]] = OpFunction
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[F32Ty]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[F32Ty]]
+; CHECK-NEXT: [[C:%.+]] = OpFunctionParameter [[F32Ty]]
+; CHECK-NEXT: OpLabel
+; CHECK-NEXT: [[R:%.+]] = OpExtInst [[F32Ty]] {{%.+}} fma [[A]] [[B]] [[C]]
+;; TODO: OpDecorate checks
+; CHECK-NEXT: OpReturnValue [[R]]
+; CHECK-NEXT: OpFunctionEnd
+define float @test_fma(float %a, float %b, float %c) {
+    %r = call float @llvm.fma.f32(float %a, float %b, float %c)
+    ret float %r
+}

diff  --git a/llvm/test/CodeGen/SPIRV/instructions/scalar-floating-point-arithmetic.ll b/llvm/test/CodeGen/SPIRV/instructions/scalar-floating-point-arithmetic.ll
new file mode 100644
index 0000000000000..8aa0f05ccfc25
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/instructions/scalar-floating-point-arithmetic.ll
@@ -0,0 +1,126 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
+
+; DISABLED-CHECK-DAG: OpName [[SCALAR_FNEG:%.+]] "scalar_fneg"
+; CHECK-DAG: OpName [[SCALAR_FADD:%.+]] "scalar_fadd"
+; CHECK-DAG: OpName [[SCALAR_FSUB:%.+]] "scalar_fsub"
+; CHECK-DAG: OpName [[SCALAR_FMUL:%.+]] "scalar_fmul"
+; CHECK-DAG: OpName [[SCALAR_FDIV:%.+]] "scalar_fdiv"
+; CHECK-DAG: OpName [[SCALAR_FREM:%.+]] "scalar_frem"
+; CHECK-DAG: OpName [[SCALAR_FMA:%.+]] "scalar_fma"
+;; FIXME: add test for OpFMod
+
+; CHECK-NOT: DAG-FENCE
+
+; CHECK-DAG: [[SCALAR:%.+]] = OpTypeFloat 32
+; CHECK-DAG: [[SCALAR_FN:%.+]] = OpTypeFunction [[SCALAR]] [[SCALAR]] [[SCALAR]]
+
+; CHECK-NOT: DAG-FENCE
+
+
+;; Test fneg on scalar:
+;; FIXME: Uncomment this test once we have rebased onto a more recent LLVM
+;;        version -- IRTranslator::translateFNeg was fixed.
+;; define float @scalar_fneg(float %a, float %unused) {
+;;     %c = fneg float %a
+;;     ret float %c
+;; }
+
+; DISABLED-CHECK:      [[SCALAR_FNEG]] = OpFunction [[SCALAR]] None [[SCALAR_FN]]
+; DISABLED-CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[SCALAR]]
+; DISABLED-CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[SCALAR]]
+; DISABLED-CHECK:      OpLabel
+; DISABLED-CHECK:      [[C:%.+]] = OpFNegate [[SCALAR]] [[A]]
+; DISABLED-CHECK:      OpReturnValue [[C]]
+; DISABLED-CHECK-NEXT: OpFunctionEnd
+
+
+;; Test fadd on scalar:
+define float @scalar_fadd(float %a, float %b) {
+    %c = fadd float %a, %b
+    ret float %c
+}
+
+; CHECK:      [[SCALAR_FADD]] = OpFunction [[SCALAR]] None [[SCALAR_FN]]
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[SCALAR]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[SCALAR]]
+; CHECK:      OpLabel
+; CHECK:      [[C:%.+]] = OpFAdd [[SCALAR]] [[A]] [[B]]
+; CHECK:      OpReturnValue [[C]]
+; CHECK-NEXT: OpFunctionEnd
+
+
+;; Test fsub on scalar:
+define float @scalar_fsub(float %a, float %b) {
+    %c = fsub float %a, %b
+    ret float %c
+}
+
+; CHECK:      [[SCALAR_FSUB]] = OpFunction [[SCALAR]] None [[SCALAR_FN]]
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[SCALAR]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[SCALAR]]
+; CHECK:      OpLabel
+; CHECK:      [[C:%.+]] = OpFSub [[SCALAR]] [[A]] [[B]]
+; CHECK:      OpReturnValue [[C]]
+; CHECK-NEXT: OpFunctionEnd
+
+
+;; Test fmul on scalar:
+define float @scalar_fmul(float %a, float %b) {
+    %c = fmul float %a, %b
+    ret float %c
+}
+
+; CHECK:      [[SCALAR_FMUL]] = OpFunction [[SCALAR]] None [[SCALAR_FN]]
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[SCALAR]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[SCALAR]]
+; CHECK:      OpLabel
+; CHECK:      [[C:%.+]] = OpFMul [[SCALAR]] [[A]] [[B]]
+; CHECK:      OpReturnValue [[C]]
+; CHECK-NEXT: OpFunctionEnd
+
+
+;; Test fdiv on scalar:
+define float @scalar_fdiv(float %a, float %b) {
+    %c = fdiv float %a, %b
+    ret float %c
+}
+
+; CHECK:      [[SCALAR_FDIV]] = OpFunction [[SCALAR]] None [[SCALAR_FN]]
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[SCALAR]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[SCALAR]]
+; CHECK:      OpLabel
+; CHECK:      [[C:%.+]] = OpFDiv [[SCALAR]] [[A]] [[B]]
+; CHECK:      OpReturnValue [[C]]
+; CHECK-NEXT: OpFunctionEnd
+
+
+;; Test frem on scalar:
+define float @scalar_frem(float %a, float %b) {
+    %c = frem float %a, %b
+    ret float %c
+}
+
+; CHECK:      [[SCALAR_FREM]] = OpFunction [[SCALAR]] None [[SCALAR_FN]]
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[SCALAR]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[SCALAR]]
+; CHECK:      OpLabel
+; CHECK:      [[C:%.+]] = OpFRem [[SCALAR]] [[A]] [[B]]
+; CHECK:      OpReturnValue [[C]]
+; CHECK-NEXT: OpFunctionEnd
+
+declare float @llvm.fma.f32(float, float, float)
+
+;; Test fma on scalar:
+define float @scalar_fma(float %a, float %b, float %c) {
+    %r = call float @llvm.fma.f32(float %a, float %b, float %c)
+    ret float %r
+}
+
+; CHECK:      [[SCALAR_FMA]] = OpFunction
+; CHECK-NEXT: [[A:%.+]] = OpFunctionParameter [[SCALAR]]
+; CHECK-NEXT: [[B:%.+]] = OpFunctionParameter [[SCALAR]]
+; CHECK-NEXT: [[C:%.+]] = OpFunctionParameter [[SCALAR]]
+; CHECK:      OpLabel
+; CHECK:      [[R:%.+]] = OpExtInst [[SCALAR]] {{%.+}} fma [[A]] [[B]] [[C]]
+; CHECK:      OpReturnValue [[R]]
+; CHECK-NEXT: OpFunctionEnd

diff  --git a/llvm/test/CodeGen/SPIRV/link-attribute.ll b/llvm/test/CodeGen/SPIRV/link-attribute.ll
new file mode 100644
index 0000000000000..172950e4ca381
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/link-attribute.ll
@@ -0,0 +1,33 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+
+%opencl.image2d_t = type opaque
+
+; CHECK: OpDecorate %[[#ID:]] LinkageAttributes "imageSampler" Export
+; CHECK: %[[#ID]] = OpVariable %[[#]] UniformConstant %[[#]]
+
+ at imageSampler = addrspace(2) constant i32 36, align 4
+
+define spir_kernel void @sample_kernel(%opencl.image2d_t addrspace(1)* %input, float addrspace(1)* nocapture %xOffsets, float addrspace(1)* nocapture %yOffsets, <4 x float> addrspace(1)* nocapture %results) {
+  %1 = tail call spir_func i64 @_Z13get_global_idj(i32 0)
+  %2 = trunc i64 %1 to i32
+  %3 = tail call spir_func i64 @_Z13get_global_idj(i32 1)
+  %4 = trunc i64 %3 to i32
+  %5 = tail call spir_func i32 @_Z15get_image_width11ocl_image2d(%opencl.image2d_t addrspace(1)* %input)
+  %6 = mul nsw i32 %4, %5
+  %7 = add nsw i32 %6, %2
+  %8 = sitofp i32 %2 to float
+  %9 = insertelement <2 x float> undef, float %8, i32 0
+  %10 = sitofp i32 %4 to float
+  %11 = insertelement <2 x float> %9, float %10, i32 1
+  %12 = tail call spir_func <4 x float> @_Z11read_imagef11ocl_image2d11ocl_samplerDv2_f(%opencl.image2d_t addrspace(1)* %input, i32 36, <2 x float> %11)
+  %13 = sext i32 %7 to i64
+  %14 = getelementptr inbounds <4 x float>, <4 x float> addrspace(1)* %results, i64 %13
+  store <4 x float> %12, <4 x float> addrspace(1)* %14, align 16
+  ret void
+}
+
+declare spir_func i64 @_Z13get_global_idj(i32)
+
+declare spir_func i32 @_Z15get_image_width11ocl_image2d(%opencl.image2d_t addrspace(1)*)
+
+declare spir_func <4 x float> @_Z11read_imagef11ocl_image2d11ocl_samplerDv2_f(%opencl.image2d_t addrspace(1)*, i32, <2 x float>)

diff  --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/abs.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/abs.ll
new file mode 100644
index 0000000000000..d627a2fc7838a
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/abs.ll
@@ -0,0 +1,21 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-linux %s -o - | FileCheck %s
+
+; CHECK: %[[#]] = OpExtInst %[[#]] %[[#]] s_abs
+; CHECK: %[[#]] = OpExtInst %[[#]] %[[#]] s_abs
+
+ at ga = addrspace(1) global i32 undef, align 4
+ at gb = addrspace(1) global <4 x i32> undef, align 4
+
+define dso_local spir_kernel void @test(i32 %a, <4 x i32> %b) local_unnamed_addr {
+entry:
+  %0 = tail call i32 @llvm.abs.i32(i32 %a, i1 0)
+  store i32 %0, i32 addrspace(1)* @ga, align 4
+  %1 = tail call <4 x i32> @llvm.abs.v4i32(<4 x i32> %b, i1 0)
+  store <4 x i32> %1, <4 x i32> addrspace(1)* @gb, align 4
+
+  ret void
+}
+
+declare i32 @llvm.abs.i32(i32, i1)
+
+declare <4 x i32> @llvm.abs.v4i32(<4 x i32>, i1)

diff  --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ceil.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ceil.ll
new file mode 100644
index 0000000000000..8411d90e23f0d
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ceil.ll
@@ -0,0 +1,43 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+
+; CHECK: %[[#extinst_id:]] = OpExtInstImport "OpenCL.std"
+
+; CHECK: %[[#var1:]] = OpTypeFloat 32
+; CHECK: %[[#var2:]] = OpTypeFloat 64
+; CHECK: %[[#var3:]] = OpTypeVector %[[#var1]] 4
+
+; CHECK: OpFunction
+; CHECK: %[[#]] = OpExtInst %[[#var1]] %[[#extinst_id]] ceil
+; CHECK: OpFunctionEnd
+
+define spir_func float @TestCeil32(float %x) local_unnamed_addr {
+entry:
+  %0 = tail call float @llvm.ceil.f32(float %x)
+  ret float %0
+}
+
+; CHECK: OpFunction
+; CHECK: %[[#]] = OpExtInst %[[#var2]] %[[#extinst_id]] ceil
+; CHECK: OpFunctionEnd
+
+define spir_func double @TestCeil64(double %x) local_unnamed_addr {
+entry:
+  %0 = tail call double @llvm.ceil.f64(double %x)
+  ret double %0
+}
+
+; CHECK: OpFunction
+; CHECK: %[[#]] = OpExtInst %[[#var3]] %[[#extinst_id]] ceil
+; CHECK: OpFunctionEnd
+
+define spir_func <4 x float> @TestCeilVec(<4 x float> %x) local_unnamed_addr {
+entry:
+  %0 = tail call <4 x float> @llvm.ceil.v4f32(<4 x float> %x)
+  ret <4 x float> %0
+}
+
+declare float @llvm.ceil.f32(float)
+
+declare double @llvm.ceil.f64(double)
+
+declare <4 x float> @llvm.ceil.v4f32(<4 x float>)

diff  --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctlz.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctlz.ll
new file mode 100644
index 0000000000000..147fd80462f9a
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/ctlz.ll
@@ -0,0 +1,15 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+
+; CHECK: %[[#extinst_id:]] = OpExtInstImport "OpenCL.std"
+
+; CHECK: OpFunction
+; CHECK: %[[#]] = OpExtInst %[[#]] %[[#extinst_id]] clz
+; CHECK: OpFunctionEnd
+
+define spir_func i32 @TestClz(i32 %x) local_unnamed_addr {
+entry:
+  %0 = tail call i32 @llvm.ctlz.i32(i32 %x, i1 true)
+  ret i32 %0
+}
+
+declare i32 @llvm.ctlz.i32(i32, i1 immarg)

diff  --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/cttz.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/cttz.ll
new file mode 100644
index 0000000000000..b9bdcfdb162f7
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/cttz.ll
@@ -0,0 +1,27 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+
+; CHECK: %[[#extinst_id:]] = OpExtInstImport "OpenCL.std"
+
+; CHECK: OpFunction
+; CHECK: %[[#]] = OpExtInst %[[#]] %[[#extinst_id]] ctz
+; CHECK: OpFunctionEnd
+
+define spir_func i32 @TestCtz(i32 %x) local_unnamed_addr {
+entry:
+  %0 = tail call i32 @llvm.cttz.i32(i32 %x, i1 true)
+  ret i32 %0
+}
+
+; CHECK: OpFunction
+; CHECK: %[[#]] = OpExtInst %[[#]] %[[#extinst_id]] ctz
+; CHECK: OpFunctionEnd
+
+define spir_func <4 x i32> @TestCtzVec(<4 x i32> %x) local_unnamed_addr {
+entry:
+  %0 = tail call <4 x i32> @llvm.cttz.v4i32(<4 x i32> %x, i1 true)
+  ret <4 x i32> %0
+}
+
+declare i32 @llvm.cttz.i32(i32, i1 immarg)
+
+declare <4 x i32> @llvm.cttz.v4i32(<4 x i32>, i1 immarg)

diff  --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/fabs.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/fabs.ll
new file mode 100644
index 0000000000000..5de2202176292
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/fabs.ll
@@ -0,0 +1,56 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+
+; CHECK: %[[#extinst_id:]] = OpExtInstImport "OpenCL.std"
+
+; CHECK: %[[#var0:]] = OpTypeFloat 16
+; CHECK: %[[#var1:]] = OpTypeFloat 32
+; CHECK: %[[#var2:]] = OpTypeFloat 64
+; CHECK: %[[#var3:]] = OpTypeVector %[[#var1]] 4
+
+; CHECK: OpFunction
+; CHECK: %[[#]] = OpExtInst %[[#var0]] %[[#extinst_id]] fabs
+; CHECK: OpFunctionEnd
+
+define spir_func half @TestFabs16(half %x) local_unnamed_addr {
+entry:
+  %0 = tail call half @llvm.fabs.f16(half %x)
+  ret half %0
+}
+
+; CHECK: OpFunction
+; CHECK: %[[#]] = OpExtInst %[[#var1]] %[[#extinst_id]] fabs
+; CHECK: OpFunctionEnd
+
+define spir_func float @TestFabs32(float %x) local_unnamed_addr {
+entry:
+  %0 = tail call float @llvm.fabs.f32(float %x)
+  ret float %0
+}
+
+; CHECK: OpFunction
+; CHECK: %[[#]] = OpExtInst %[[#var2]] %[[#extinst_id]] fabs
+; CHECK: OpFunctionEnd
+
+define spir_func double @TestFabs64(double %x) local_unnamed_addr {
+entry:
+  %0 = tail call double @llvm.fabs.f64(double %x)
+  ret double %0
+}
+
+; CHECK: OpFunction
+; CHECK: %[[#]] = OpExtInst %[[#var3]] %[[#extinst_id]] fabs
+; CHECK: OpFunctionEnd
+
+define spir_func <4 x float> @TestFabsVec(<4 x float> %x) local_unnamed_addr {
+entry:
+  %0 = tail call <4 x float> @llvm.fabs.v4f32(<4 x float> %x)
+  ret <4 x float> %0
+}
+
+declare half @llvm.fabs.f16(half)
+
+declare float @llvm.fabs.f32(float)
+
+declare double @llvm.fabs.f64(double)
+
+declare <4 x float> @llvm.fabs.v4f32(<4 x float>)

diff  --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/nearbyint.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/nearbyint.ll
new file mode 100644
index 0000000000000..70194e21982dc
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/nearbyint.ll
@@ -0,0 +1,11 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
+
+; CHECK: %[[#]] = OpExtInst %[[#]] %[[#]] rint
+
+define dso_local spir_func float @foo(float %x) local_unnamed_addr {
+entry:
+  %0 = tail call float @llvm.nearbyint.f32(float %x)
+  ret float %0
+}
+
+declare float @llvm.nearbyint.f32(float)

diff  --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/sqrt.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/sqrt.ll
new file mode 100644
index 0000000000000..4b1d7535dba32
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/sqrt.ll
@@ -0,0 +1,30 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
+
+; CHECK: %[[#ExtInstSetId:]] = OpExtInstImport "OpenCL.std"
+; CHECK: %[[#Float:]] = OpTypeFloat 32
+; CHECK: %[[#Double:]] = OpTypeFloat 64
+; CHECK: %[[#Double4:]] = OpTypeVector %[[#Double]] 4
+; CHECK: %[[#FloatArg:]] = OpConstant %[[#Float]]
+; CHECK: %[[#DoubleArg:]] = OpConstant %[[#Double]]
+; CHECK: %[[#Double4Arg:]] = OpConstantComposite %[[#Double4]]
+
+;; We need to store sqrt results, otherwise isel does not emit sqrts as dead insts.
+define spir_func void @test_sqrt(float* %x, double* %y, <4 x double>* %z) {
+entry:
+  %0 = call float @llvm.sqrt.f32(float 0x40091EB860000000)
+  store float %0, float* %x
+  %1 = call double @llvm.sqrt.f64(double 2.710000e+00)
+  store double %1, double* %y
+  %2 = call <4 x double> @llvm.sqrt.v4f64(<4 x double> <double 5.000000e-01, double 2.000000e-01, double 3.000000e-01, double 4.000000e-01>)
+  store <4 x double> %2, <4 x double>* %z
+; CHECK: %[[#]] = OpExtInst %[[#Float]] %[[#ExtInstSetId]] sqrt %[[#FloatArg]]
+; CHECK: %[[#]] = OpExtInst %[[#Double]] %[[#ExtInstSetId]] sqrt %[[#DoubleArg]]
+; CHECK: %[[#]] = OpExtInst %[[#Double4]] %[[#ExtInstSetId]] sqrt %[[#Double4Arg]]
+  ret void
+}
+
+declare float @llvm.sqrt.f32(float)
+
+declare double @llvm.sqrt.f64(double)
+
+declare <4 x double> @llvm.sqrt.v4f64(<4 x double>)

diff  --git a/llvm/test/CodeGen/SPIRV/mangled_function.ll b/llvm/test/CodeGen/SPIRV/mangled_function.ll
new file mode 100644
index 0000000000000..4185d68290ca5
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/mangled_function.ll
@@ -0,0 +1,25 @@
+;; Make sure backend doesn't crash if the program contains
+;; a mangled function which is not an OpenCL bultin.
+;; Source:
+;; void __attribute__((overloadable))
+;; foo(image2d_t srcImage);
+;;
+;; void bar(image2d_t srcImage) {
+;;   foo(srcImage);
+;; }
+;; clang -cc1 /work/tmp/tmp.cl -cl-std=CL2.0 -triple spir-unknown-unknown  -finclude-default-header -emit-llvm -o test/mangled_function.ll
+
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: OpName %[[#foo:]] "_Z3foo14ocl_image2d_ro"
+; CHECK-SPIRV: %[[#foo]] = OpFunction %[[#]]
+
+%opencl.image2d_ro_t = type opaque
+
+define spir_func void @bar(%opencl.image2d_ro_t addrspace(1)* %srcImage) local_unnamed_addr {
+entry:
+  tail call spir_func void @_Z3foo14ocl_image2d_ro(%opencl.image2d_ro_t addrspace(1)* %srcImage)
+  ret void
+}
+
+declare spir_func void @_Z3foo14ocl_image2d_ro(%opencl.image2d_ro_t addrspace(1)*) local_unnamed_addr

diff  --git a/llvm/test/CodeGen/SPIRV/no_capability_shader.ll b/llvm/test/CodeGen/SPIRV/no_capability_shader.ll
new file mode 100644
index 0000000000000..4121c03e7a830
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/no_capability_shader.ll
@@ -0,0 +1,13 @@
+;; __kernel void sample_test(read_only image2d_t src, read_only image1d_buffer_t buff) {}
+
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV-NOT: OpCapability Shader
+
+%opencl.image2d_ro_t = type opaque
+%opencl.image1d_buffer_ro_t = type opaque
+
+define spir_kernel void @sample_test(%opencl.image2d_ro_t addrspace(1)* %src, %opencl.image1d_buffer_ro_t addrspace(1)* %buf) {
+entry:
+  ret void
+}

diff  --git a/llvm/test/CodeGen/SPIRV/opencl.queue_t.ll b/llvm/test/CodeGen/SPIRV/opencl.queue_t.ll
new file mode 100644
index 0000000000000..5c07fe6e3e3f5
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/opencl.queue_t.ll
@@ -0,0 +1,11 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: OpCapability DeviceEnqueue
+; CHECK-SPIRV: OpTypeQueue
+
+%opencl.queue_t = type opaque
+
+define spir_func void @enqueue_simple_block(%opencl.queue_t* addrspace(3)* nocapture %q) {
+entry:
+  ret void
+}

diff  --git a/llvm/test/CodeGen/SPIRV/read_image.ll b/llvm/test/CodeGen/SPIRV/read_image.ll
new file mode 100644
index 0000000000000..1fdcf9b5d9086
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/read_image.ll
@@ -0,0 +1,50 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: %[[#IntTy:]] = OpTypeInt
+; CHECK-SPIRV: %[[#IVecTy:]] = OpTypeVector %[[#IntTy]]
+; CHECK-SPIRV: %[[#FloatTy:]] = OpTypeFloat
+; CHECK-SPIRV: %[[#FVecTy:]] = OpTypeVector %[[#FloatTy]]
+; CHECK-SPIRV: OpImageRead %[[#IVecTy]]
+; CHECK-SPIRV: OpImageRead %[[#FVecTy]]
+
+;; __kernel void kernelA(__read_only image3d_t input) {
+;;   uint4 c = read_imageui(input, (int4)(0, 0, 0, 0));
+;; }
+;;
+;; __kernel void kernelB(__read_only image3d_t input) {
+;;   float4 f = read_imagef(input, (int4)(0, 0, 0, 0));
+;; }
+
+%opencl.image3d_ro_t = type opaque
+
+define dso_local spir_kernel void @kernelA(%opencl.image3d_ro_t addrspace(1)* %input) {
+entry:
+  %input.addr = alloca %opencl.image3d_ro_t addrspace(1)*, align 8
+  %c = alloca <4 x i32>, align 16
+  %.compoundliteral = alloca <4 x i32>, align 16
+  store %opencl.image3d_ro_t addrspace(1)* %input, %opencl.image3d_ro_t addrspace(1)** %input.addr, align 8
+  %0 = load %opencl.image3d_ro_t addrspace(1)*, %opencl.image3d_ro_t addrspace(1)** %input.addr, align 8
+  store <4 x i32> zeroinitializer, <4 x i32>* %.compoundliteral, align 16
+  %1 = load <4 x i32>, <4 x i32>* %.compoundliteral, align 16
+  %call = call spir_func <4 x i32> @_Z12read_imageui14ocl_image3d_roDv4_i(%opencl.image3d_ro_t addrspace(1)* %0, <4 x i32> noundef %1)
+  store <4 x i32> %call, <4 x i32>* %c, align 16
+  ret void
+}
+
+declare spir_func <4 x i32> @_Z12read_imageui14ocl_image3d_roDv4_i(%opencl.image3d_ro_t addrspace(1)*, <4 x i32> noundef)
+
+define dso_local spir_kernel void @kernelB(%opencl.image3d_ro_t addrspace(1)* %input) {
+entry:
+  %input.addr = alloca %opencl.image3d_ro_t addrspace(1)*, align 8
+  %f = alloca <4 x float>, align 16
+  %.compoundliteral = alloca <4 x i32>, align 16
+  store %opencl.image3d_ro_t addrspace(1)* %input, %opencl.image3d_ro_t addrspace(1)** %input.addr, align 8
+  %0 = load %opencl.image3d_ro_t addrspace(1)*, %opencl.image3d_ro_t addrspace(1)** %input.addr, align 8
+  store <4 x i32> zeroinitializer, <4 x i32>* %.compoundliteral, align 16
+  %1 = load <4 x i32>, <4 x i32>* %.compoundliteral, align 16
+  %call = call spir_func <4 x float> @_Z11read_imagef14ocl_image3d_roDv4_i(%opencl.image3d_ro_t addrspace(1)* %0, <4 x i32> noundef %1)
+  store <4 x float> %call, <4 x float>* %f, align 16
+  ret void
+}
+
+declare spir_func <4 x float> @_Z11read_imagef14ocl_image3d_roDv4_i(%opencl.image3d_ro_t addrspace(1)*, <4 x i32> noundef)

diff  --git a/llvm/test/CodeGen/SPIRV/spirv.Queue.ll b/llvm/test/CodeGen/SPIRV/spirv.Queue.ll
new file mode 100644
index 0000000000000..6c511477cfa8a
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/spirv.Queue.ll
@@ -0,0 +1,11 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: OpCapability DeviceEnqueue
+; CHECK-SPIRV: OpTypeQueue
+
+%spirv.Queue = type opaque
+
+define spir_func void @enqueue_simple_block(%spirv.Queue* addrspace(3)* nocapture %q) {
+entry:
+  ret void
+}

diff  --git a/llvm/test/CodeGen/SPIRV/transcoding/OpConstantSampler.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpConstantSampler.ll
new file mode 100644
index 0000000000000..c85a9c76c5055
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpConstantSampler.ll
@@ -0,0 +1,35 @@
+;; Sources:
+;;
+;; void kernel foo(__read_only image2d_t src) {
+;;   sampler_t sampler1 = CLK_NORMALIZED_COORDS_TRUE |
+;;                        CLK_ADDRESS_REPEAT |
+;;                        CLK_FILTER_NEAREST;
+;;   sampler_t sampler2 = 0x00;
+;;
+;;   read_imagef(src, sampler1, 0, 0);
+;;   read_imagef(src, sampler2, 0, 0);
+;; }
+
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: %[[#SamplerID0:]] = OpConstantSampler %[[#]] Repeat 1 Nearest
+; CHECK-SPIRV: %[[#SamplerID1:]] = OpConstantSampler %[[#]] None 0 Nearest
+; CHECK-SPIRV: %[[#]] = OpSampledImage %[[#]] %[[#]] %[[#SamplerID0]]
+; CHECK-SPIRV: %[[#]] = OpSampledImage %[[#]] %[[#]] %[[#SamplerID1]]
+
+%opencl.image2d_ro_t = type opaque
+%opencl.sampler_t = type opaque
+
+define spir_func <4 x float> @foo(%opencl.image2d_ro_t addrspace(1)* %src) local_unnamed_addr {
+entry:
+  %0 = tail call %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 23)
+  %1 = tail call %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 0)
+  %call = tail call spir_func <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv2_ff(%opencl.image2d_ro_t addrspace(1)* %src, %opencl.sampler_t addrspace(2)* %0, <2 x float> zeroinitializer, float 0.000000e+00)
+  %call1 = tail call spir_func <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv2_ff(%opencl.image2d_ro_t addrspace(1)* %src, %opencl.sampler_t addrspace(2)* %1, <2 x float> zeroinitializer, float 0.000000e+00)
+  %add = fadd <4 x float> %call, %call1
+  ret <4 x float> %add
+}
+
+declare %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32) local_unnamed_addr
+
+declare spir_func <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv2_ff(%opencl.image2d_ro_t addrspace(1)*, %opencl.sampler_t addrspace(2)*, <2 x float>, float) local_unnamed_addr

diff  --git a/llvm/test/CodeGen/SPIRV/transcoding/OpGroupAsyncCopy.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpGroupAsyncCopy.ll
new file mode 100644
index 0000000000000..d2d5483ed3618
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpGroupAsyncCopy.ll
@@ -0,0 +1,110 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV-DAG: %[[#]] = OpGroupAsyncCopy %[[#]] %[[#Scope:]]
+; CHECK-SPIRV-DAG: %[[#Scope]] = OpConstant %[[#]]
+
+%opencl.event_t = type opaque
+
+define spir_kernel void @test_fn(<2 x i8> addrspace(1)* %src, <2 x i8> addrspace(1)* %dst, <2 x i8> addrspace(3)* %localBuffer, i32 %copiesPerWorkgroup, i32 %copiesPerWorkItem) {
+entry:
+  %src.addr = alloca <2 x i8> addrspace(1)*, align 4
+  %dst.addr = alloca <2 x i8> addrspace(1)*, align 4
+  %localBuffer.addr = alloca <2 x i8> addrspace(3)*, align 4
+  %copiesPerWorkgroup.addr = alloca i32, align 4
+  %copiesPerWorkItem.addr = alloca i32, align 4
+  %i = alloca i32, align 4
+  %event = alloca %opencl.event_t*, align 4
+  store <2 x i8> addrspace(1)* %src, <2 x i8> addrspace(1)** %src.addr, align 4
+  store <2 x i8> addrspace(1)* %dst, <2 x i8> addrspace(1)** %dst.addr, align 4
+  store <2 x i8> addrspace(3)* %localBuffer, <2 x i8> addrspace(3)** %localBuffer.addr, align 4
+  store i32 %copiesPerWorkgroup, i32* %copiesPerWorkgroup.addr, align 4
+  store i32 %copiesPerWorkItem, i32* %copiesPerWorkItem.addr, align 4
+  store i32 0, i32* %i, align 4
+  br label %for.cond
+
+for.cond:                                         ; preds = %for.inc, %entry
+  %0 = load i32, i32* %i, align 4
+  %1 = load i32, i32* %copiesPerWorkItem.addr, align 4
+  %cmp = icmp slt i32 %0, %1
+  br i1 %cmp, label %for.body, label %for.end
+
+for.body:                                         ; preds = %for.cond
+  %call = call spir_func i32 @_Z12get_local_idj(i32 0)
+  %2 = load i32, i32* %copiesPerWorkItem.addr, align 4
+  %mul = mul i32 %call, %2
+  %3 = load i32, i32* %i, align 4
+  %add = add i32 %mul, %3
+  %4 = load <2 x i8> addrspace(3)*, <2 x i8> addrspace(3)** %localBuffer.addr, align 4
+  %arrayidx = getelementptr inbounds <2 x i8>, <2 x i8> addrspace(3)* %4, i32 %add
+  store <2 x i8> zeroinitializer, <2 x i8> addrspace(3)* %arrayidx, align 2
+  br label %for.inc
+
+for.inc:                                          ; preds = %for.body
+  %5 = load i32, i32* %i, align 4
+  %inc = add nsw i32 %5, 1
+  store i32 %inc, i32* %i, align 4
+  br label %for.cond
+
+for.end:                                          ; preds = %for.cond
+  call spir_func void @_Z7barrierj(i32 1)
+  store i32 0, i32* %i, align 4
+  br label %for.cond1
+
+for.cond1:                                        ; preds = %for.inc12, %for.end
+  %6 = load i32, i32* %i, align 4
+  %7 = load i32, i32* %copiesPerWorkItem.addr, align 4
+  %cmp2 = icmp slt i32 %6, %7
+  br i1 %cmp2, label %for.body3, label %for.end14
+
+for.body3:                                        ; preds = %for.cond1
+  %call4 = call spir_func i32 @_Z13get_global_idj(i32 0)
+  %8 = load i32, i32* %copiesPerWorkItem.addr, align 4
+  %mul5 = mul i32 %call4, %8
+  %9 = load i32, i32* %i, align 4
+  %add6 = add i32 %mul5, %9
+  %10 = load <2 x i8> addrspace(1)*, <2 x i8> addrspace(1)** %src.addr, align 4
+  %arrayidx7 = getelementptr inbounds <2 x i8>, <2 x i8> addrspace(1)* %10, i32 %add6
+  %11 = load <2 x i8>, <2 x i8> addrspace(1)* %arrayidx7, align 2
+  %call8 = call spir_func i32 @_Z12get_local_idj(i32 0)
+  %12 = load i32, i32* %copiesPerWorkItem.addr, align 4
+  %mul9 = mul i32 %call8, %12
+  %13 = load i32, i32* %i, align 4
+  %add10 = add i32 %mul9, %13
+  %14 = load <2 x i8> addrspace(3)*, <2 x i8> addrspace(3)** %localBuffer.addr, align 4
+  %arrayidx11 = getelementptr inbounds <2 x i8>, <2 x i8> addrspace(3)* %14, i32 %add10
+  store <2 x i8> %11, <2 x i8> addrspace(3)* %arrayidx11, align 2
+  br label %for.inc12
+
+for.inc12:                                        ; preds = %for.body3
+  %15 = load i32, i32* %i, align 4
+  %inc13 = add nsw i32 %15, 1
+  store i32 %inc13, i32* %i, align 4
+  br label %for.cond1
+
+for.end14:                                        ; preds = %for.cond1
+  call spir_func void @_Z7barrierj(i32 1)
+  %16 = load <2 x i8> addrspace(1)*, <2 x i8> addrspace(1)** %dst.addr, align 4
+  %17 = load i32, i32* %copiesPerWorkgroup.addr, align 4
+  %call15 = call spir_func i32 @_Z12get_group_idj(i32 0)
+  %mul16 = mul i32 %17, %call15
+  %add.ptr = getelementptr inbounds <2 x i8>, <2 x i8> addrspace(1)* %16, i32 %mul16
+  %18 = load <2 x i8> addrspace(3)*, <2 x i8> addrspace(3)** %localBuffer.addr, align 4
+  %19 = load i32, i32* %copiesPerWorkgroup.addr, align 4
+  %call17 = call spir_func %opencl.event_t* @_Z21async_work_group_copyPU3AS1Dv2_cPKU3AS3S_j9ocl_event(<2 x i8> addrspace(1)* %add.ptr, <2 x i8> addrspace(3)* %18, i32 %19, %opencl.event_t* null)
+  store %opencl.event_t* %call17, %opencl.event_t** %event, align 4
+  %20 = addrspacecast %opencl.event_t** %event to %opencl.event_t* addrspace(4)*
+  call spir_func void @_Z17wait_group_eventsiPU3AS49ocl_event(i32 1, %opencl.event_t* addrspace(4)* %20)
+  ret void
+}
+
+declare spir_func i32 @_Z12get_local_idj(i32)
+
+declare spir_func void @_Z7barrierj(i32)
+
+declare spir_func i32 @_Z13get_global_idj(i32)
+
+declare spir_func %opencl.event_t* @_Z21async_work_group_copyPU3AS1Dv2_cPKU3AS3S_j9ocl_event(<2 x i8> addrspace(1)*, <2 x i8> addrspace(3)*, i32, %opencl.event_t*)
+
+declare spir_func i32 @_Z12get_group_idj(i32)
+
+declare spir_func void @_Z17wait_group_eventsiPU3AS49ocl_event(i32, %opencl.event_t* addrspace(4)*)

diff  --git a/llvm/test/CodeGen/SPIRV/transcoding/OpImageQuerySize.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpImageQuerySize.ll
new file mode 100644
index 0000000000000..ca80dcd6999cd
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpImageQuerySize.ll
@@ -0,0 +1,126 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+;; Check conversion of get_image_width, get_image_height, get_image_depth,
+;; get_image_array_size, and get_image_dim OCL built-ins.
+;; In general the SPRI-V reader converts OpImageQuerySize into get_image_dim
+;; and subsequent extract or shufflevector instructions. Unfortunately there is
+;; no get_image_dim for 1D images and get_image_dim cannot replace get_image_array_size
+
+; CHECK-SPIRV: %[[#ArrayTypeID:]] = OpTypeImage %[[#]] 1D 0 1 0 0 Unknown ReadOnly
+
+%opencl.image1d_ro_t = type opaque
+%opencl.image1d_buffer_ro_t = type opaque
+%opencl.image1d_array_ro_t = type opaque
+%opencl.image2d_ro_t = type opaque
+%opencl.image2d_depth_ro_t = type opaque
+%opencl.image2d_array_ro_t = type opaque
+%opencl.image2d_array_depth_ro_t = type opaque
+%opencl.image3d_ro_t = type opaque
+
+; CHECK-SPIRV:     %[[#ArrayVarID:]] = OpFunctionParameter %[[#ArrayTypeID]]
+; CHECK-SPIRV:     %[[#]] = OpImageQuerySizeLod %[[#]] %[[#ArrayVarID]]
+; CHECK-SPIRV-NOT: %[[#]] = OpExtInst %[[#]] %[[#]] get_image_array_size
+
+define spir_kernel void @test_image1d(i32 addrspace(1)* nocapture %sizes, %opencl.image1d_ro_t addrspace(1)* %img, %opencl.image1d_buffer_ro_t addrspace(1)* %buffer, %opencl.image1d_array_ro_t addrspace(1)* %array) {
+  %1 = tail call spir_func i32 @_Z15get_image_width14ocl_image1d_ro(%opencl.image1d_ro_t addrspace(1)* %img)
+  %2 = tail call spir_func i32 @_Z15get_image_width21ocl_image1d_buffer_ro(%opencl.image1d_buffer_ro_t addrspace(1)* %buffer)
+  %3 = tail call spir_func i32 @_Z15get_image_width20ocl_image1d_array_ro(%opencl.image1d_array_ro_t addrspace(1)* %array)
+  %4 = tail call spir_func i64 @_Z20get_image_array_size20ocl_image1d_array_ro(%opencl.image1d_array_ro_t addrspace(1)* %array)
+  %5 = trunc i64 %4 to i32
+  %6 = add nsw i32 %2, %1
+  %7 = add nsw i32 %6, %3
+  %8 = add nsw i32 %7, %5
+  store i32 %8, i32 addrspace(1)* %sizes, align 4
+  ret void
+}
+
+declare spir_func i32 @_Z15get_image_width14ocl_image1d_ro(%opencl.image1d_ro_t addrspace(1)*)
+
+declare spir_func i32 @_Z15get_image_width21ocl_image1d_buffer_ro(%opencl.image1d_buffer_ro_t addrspace(1)*)
+
+declare spir_func i32 @_Z15get_image_width20ocl_image1d_array_ro(%opencl.image1d_array_ro_t addrspace(1)*)
+
+declare spir_func i64 @_Z20get_image_array_size20ocl_image1d_array_ro(%opencl.image1d_array_ro_t addrspace(1)*)
+
+define spir_kernel void @test_image2d(i32 addrspace(1)* nocapture %sizes, %opencl.image2d_ro_t addrspace(1)* %img, %opencl.image2d_depth_ro_t addrspace(1)* nocapture %img_depth, %opencl.image2d_array_ro_t addrspace(1)* %array, %opencl.image2d_array_depth_ro_t addrspace(1)* nocapture %array_depth) {
+  %1 = tail call spir_func i32 @_Z15get_image_width14ocl_image2d_ro(%opencl.image2d_ro_t addrspace(1)* %img)
+  %2 = tail call spir_func i32 @_Z16get_image_height14ocl_image2d_ro(%opencl.image2d_ro_t addrspace(1)* %img)
+  %3 = tail call spir_func <2 x i32> @_Z13get_image_dim14ocl_image2d_ro(%opencl.image2d_ro_t addrspace(1)* %img)
+  %4 = tail call spir_func i32 @_Z15get_image_width20ocl_image2d_array_ro(%opencl.image2d_array_ro_t addrspace(1)* %array)
+  %5 = tail call spir_func i32 @_Z16get_image_height20ocl_image2d_array_ro(%opencl.image2d_array_ro_t addrspace(1)* %array)
+  %6 = tail call spir_func i64 @_Z20get_image_array_size20ocl_image2d_array_ro(%opencl.image2d_array_ro_t addrspace(1)* %array)
+  %7 = trunc i64 %6 to i32
+  %8 = tail call spir_func <2 x i32> @_Z13get_image_dim20ocl_image2d_array_ro(%opencl.image2d_array_ro_t addrspace(1)* %array)
+  %9 = add nsw i32 %2, %1
+  %10 = extractelement <2 x i32> %3, i32 0
+  %11 = add nsw i32 %9, %10
+  %12 = extractelement <2 x i32> %3, i32 1
+  %13 = add nsw i32 %11, %12
+  %14 = add nsw i32 %13, %4
+  %15 = add nsw i32 %14, %5
+  %16 = add nsw i32 %15, %7
+  %17 = extractelement <2 x i32> %8, i32 0
+  %18 = add nsw i32 %16, %17
+  %19 = extractelement <2 x i32> %8, i32 1
+  %20 = add nsw i32 %18, %19
+  store i32 %20, i32 addrspace(1)* %sizes, align 4
+  ret void
+}
+
+declare spir_func i32 @_Z15get_image_width14ocl_image2d_ro(%opencl.image2d_ro_t addrspace(1)*)
+
+declare spir_func i32 @_Z16get_image_height14ocl_image2d_ro(%opencl.image2d_ro_t addrspace(1)*)
+
+declare spir_func <2 x i32> @_Z13get_image_dim14ocl_image2d_ro(%opencl.image2d_ro_t addrspace(1)*)
+
+declare spir_func i32 @_Z15get_image_width20ocl_image2d_array_ro(%opencl.image2d_array_ro_t addrspace(1)*)
+
+declare spir_func i32 @_Z16get_image_height20ocl_image2d_array_ro(%opencl.image2d_array_ro_t addrspace(1)*)
+
+declare spir_func i64 @_Z20get_image_array_size20ocl_image2d_array_ro(%opencl.image2d_array_ro_t addrspace(1)*)
+
+declare spir_func <2 x i32> @_Z13get_image_dim20ocl_image2d_array_ro(%opencl.image2d_array_ro_t addrspace(1)*)
+
+define spir_kernel void @test_image3d(i32 addrspace(1)* nocapture %sizes, %opencl.image3d_ro_t addrspace(1)* %img) {
+  %1 = tail call spir_func i32 @_Z15get_image_width14ocl_image3d_ro(%opencl.image3d_ro_t addrspace(1)* %img)
+  %2 = tail call spir_func i32 @_Z16get_image_height14ocl_image3d_ro(%opencl.image3d_ro_t addrspace(1)* %img)
+  %3 = tail call spir_func i32 @_Z15get_image_depth14ocl_image3d_ro(%opencl.image3d_ro_t addrspace(1)* %img)
+  %4 = tail call spir_func <4 x i32> @_Z13get_image_dim14ocl_image3d_ro(%opencl.image3d_ro_t addrspace(1)* %img)
+  %5 = add nsw i32 %2, %1
+  %6 = add nsw i32 %5, %3
+  %7 = extractelement <4 x i32> %4, i32 0
+  %8 = add nsw i32 %6, %7
+  %9 = extractelement <4 x i32> %4, i32 1
+  %10 = add nsw i32 %8, %9
+  %11 = extractelement <4 x i32> %4, i32 2
+  %12 = add nsw i32 %10, %11
+  %13 = extractelement <4 x i32> %4, i32 3
+  %14 = add nsw i32 %12, %13
+  store i32 %14, i32 addrspace(1)* %sizes, align 4
+  ret void
+}
+
+declare spir_func i32 @_Z15get_image_width14ocl_image3d_ro(%opencl.image3d_ro_t addrspace(1)*)
+
+declare spir_func i32 @_Z16get_image_height14ocl_image3d_ro(%opencl.image3d_ro_t addrspace(1)*)
+
+declare spir_func i32 @_Z15get_image_depth14ocl_image3d_ro(%opencl.image3d_ro_t addrspace(1)*)
+
+declare spir_func <4 x i32> @_Z13get_image_dim14ocl_image3d_ro(%opencl.image3d_ro_t addrspace(1)*)
+
+define spir_kernel void @test_image2d_array_depth_t(i32 addrspace(1)* nocapture %sizes, %opencl.image2d_array_depth_ro_t addrspace(1)* %array) {
+  %1 = tail call spir_func i32 @_Z15get_image_width26ocl_image2d_array_depth_ro(%opencl.image2d_array_depth_ro_t addrspace(1)* %array)
+  %2 = tail call spir_func i32 @_Z16get_image_height26ocl_image2d_array_depth_ro(%opencl.image2d_array_depth_ro_t addrspace(1)* %array)
+  %3 = tail call spir_func i64 @_Z20get_image_array_size26ocl_image2d_array_depth_ro(%opencl.image2d_array_depth_ro_t addrspace(1)* %array)
+  %4 = trunc i64 %3 to i32
+  %5 = add nsw i32 %2, %1
+  %6 = add nsw i32 %5, %4
+  store i32 %5, i32 addrspace(1)* %sizes, align 4
+  ret void
+}
+
+declare spir_func i32 @_Z15get_image_width26ocl_image2d_array_depth_ro(%opencl.image2d_array_depth_ro_t addrspace(1)*)
+
+declare spir_func i32 @_Z16get_image_height26ocl_image2d_array_depth_ro(%opencl.image2d_array_depth_ro_t addrspace(1)*)
+
+declare spir_func i64 @_Z20get_image_array_size26ocl_image2d_array_depth_ro(%opencl.image2d_array_depth_ro_t addrspace(1)*)

diff  --git a/llvm/test/CodeGen/SPIRV/transcoding/OpImageReadMS.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpImageReadMS.ll
new file mode 100644
index 0000000000000..f523c037059bf
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpImageReadMS.ll
@@ -0,0 +1,50 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: %[[#]] = OpImageRead %[[#]] %[[#]] %[[#]] Sample %[[#]]
+
+%opencl.image2d_msaa_ro_t = type opaque
+
+define spir_kernel void @sample_test(%opencl.image2d_msaa_ro_t addrspace(1)* %source, i32 %sampler, <4 x float> addrspace(1)* nocapture %results) {
+entry:
+  %call = tail call spir_func i32 @_Z13get_global_idj(i32 0)
+  %call1 = tail call spir_func i32 @_Z13get_global_idj(i32 1)
+  %call2 = tail call spir_func i32 @_Z15get_image_width19ocl_image2d_msaa_ro(%opencl.image2d_msaa_ro_t addrspace(1)* %source)
+  %call3 = tail call spir_func i32 @_Z16get_image_height19ocl_image2d_msaa_ro(%opencl.image2d_msaa_ro_t addrspace(1)* %source)
+  %call4 = tail call spir_func i32 @_Z21get_image_num_samples19ocl_image2d_msaa_ro(%opencl.image2d_msaa_ro_t addrspace(1)* %source)
+  %cmp20 = icmp eq i32 %call4, 0
+  br i1 %cmp20, label %for.end, label %for.body.lr.ph
+
+for.body.lr.ph:                                   ; preds = %entry
+  %vecinit = insertelement <2 x i32> undef, i32 %call, i32 0
+  %vecinit8 = insertelement <2 x i32> %vecinit, i32 %call1, i32 1
+  br label %for.body
+
+for.body:                                         ; preds = %for.body.lr.ph, %for.body
+  %sample.021 = phi i32 [ 0, %for.body.lr.ph ], [ %inc, %for.body ]
+  %mul5 = mul i32 %sample.021, %call3
+  %tmp = add i32 %mul5, %call1
+  %tmp19 = mul i32 %tmp, %call2
+  %add7 = add i32 %tmp19, %call
+  %call9 = tail call spir_func <4 x float> @_Z11read_imagef19ocl_image2d_msaa_roDv2_ii(%opencl.image2d_msaa_ro_t addrspace(1)* %source, <2 x i32> %vecinit8, i32 %sample.021)
+  %arrayidx = getelementptr inbounds <4 x float>, <4 x float> addrspace(1)* %results, i32 %add7
+  store <4 x float> %call9, <4 x float> addrspace(1)* %arrayidx, align 16
+  %inc = add nuw i32 %sample.021, 1
+  %cmp = icmp ult i32 %inc, %call4
+  br i1 %cmp, label %for.body, label %for.end.loopexit
+
+for.end.loopexit:                                 ; preds = %for.body
+  br label %for.end
+
+for.end:                                          ; preds = %for.end.loopexit, %entry
+  ret void
+}
+
+declare spir_func i32 @_Z13get_global_idj(i32)
+
+declare spir_func i32 @_Z15get_image_width19ocl_image2d_msaa_ro(%opencl.image2d_msaa_ro_t addrspace(1)*)
+
+declare spir_func i32 @_Z16get_image_height19ocl_image2d_msaa_ro(%opencl.image2d_msaa_ro_t addrspace(1)*)
+
+declare spir_func i32 @_Z21get_image_num_samples19ocl_image2d_msaa_ro(%opencl.image2d_msaa_ro_t addrspace(1)*)
+
+declare spir_func <4 x float> @_Z11read_imagef19ocl_image2d_msaa_roDv2_ii(%opencl.image2d_msaa_ro_t addrspace(1)*, <2 x i32>, i32)

diff  --git a/llvm/test/CodeGen/SPIRV/transcoding/OpImageSampleExplicitLod.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpImageSampleExplicitLod.ll
new file mode 100644
index 0000000000000..9c4ecad750f04
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpImageSampleExplicitLod.ll
@@ -0,0 +1,35 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV-DAG: %[[#RetID:]] = OpImageSampleExplicitLod %[[#RetType:]] %[[#]] %[[#]] Lod %[[#]]
+; CHECK-SPIRV-DAG: %[[#RetType]] = OpTypeVector %[[#]] 4
+; CHECK-SPIRV:     %[[#]] = OpCompositeExtract %[[#]] %[[#RetID]] 0
+
+%opencl.image2d_depth_ro_t = type opaque
+
+define spir_kernel void @sample_kernel(%opencl.image2d_depth_ro_t addrspace(1)* %input, i32 %imageSampler, float addrspace(1)* %xOffsets, float addrspace(1)* %yOffsets, float addrspace(1)* %results) {
+entry:
+  %call = call spir_func i32 @_Z13get_global_idj(i32 0)
+  %call1 = call spir_func i32 @_Z13get_global_idj(i32 1)
+  %call2.tmp1 = call spir_func <2 x i32> @_Z13get_image_dim20ocl_image2d_depth_ro(%opencl.image2d_depth_ro_t addrspace(1)* %input)
+  %call2.old = extractelement <2 x i32> %call2.tmp1, i32 0
+  %mul = mul i32 %call1, %call2.old
+  %add = add i32 %mul, %call
+  %arrayidx = getelementptr inbounds float, float addrspace(1)* %xOffsets, i32 %add
+  %0 = load float, float addrspace(1)* %arrayidx, align 4
+  %conv = fptosi float %0 to i32
+  %vecinit = insertelement <2 x i32> undef, i32 %conv, i32 0
+  %arrayidx3 = getelementptr inbounds float, float addrspace(1)* %yOffsets, i32 %add
+  %1 = load float, float addrspace(1)* %arrayidx3, align 4
+  %conv4 = fptosi float %1 to i32
+  %vecinit5 = insertelement <2 x i32> %vecinit, i32 %conv4, i32 1
+  %call6.tmp.tmp = call spir_func float @_Z11read_imagef20ocl_image2d_depth_ro11ocl_samplerDv2_i(%opencl.image2d_depth_ro_t addrspace(1)* %input, i32 %imageSampler, <2 x i32> %vecinit5)
+  %arrayidx7 = getelementptr inbounds float, float addrspace(1)* %results, i32 %add
+  store float %call6.tmp.tmp, float addrspace(1)* %arrayidx7, align 4
+  ret void
+}
+
+declare spir_func float @_Z11read_imagef20ocl_image2d_depth_ro11ocl_samplerDv2_i(%opencl.image2d_depth_ro_t addrspace(1)*, i32, <2 x i32>)
+
+declare spir_func i32 @_Z13get_global_idj(i32)
+
+declare spir_func <2 x i32> @_Z13get_image_dim20ocl_image2d_depth_ro(%opencl.image2d_depth_ro_t addrspace(1)*)

diff  --git a/llvm/test/CodeGen/SPIRV/transcoding/OpImageWrite.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpImageWrite.ll
new file mode 100644
index 0000000000000..7c8efa405c707
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpImageWrite.ll
@@ -0,0 +1,344 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: %[[#VOID_TY:]] = OpTypeVoid
+; CHECK-SPIRV: %[[#IMG2D_WO_TY:]] = OpTypeImage %[[#VOID_TY]] 2D 0 0 0 0 Unknown WriteOnly
+; CHECK-SPIRV: %[[#IMG2D_RW_TY:]] = OpTypeImage %[[#VOID_TY]] 2D 0 0 0 0 Unknown ReadWrite
+; CHECK-SPIRV: %[[#IMG2D_ARRAY_WO_TY:]] = OpTypeImage %[[#VOID_TY]] 2D 0 1 0 0 Unknown WriteOnly
+; CHECK-SPIRV: %[[#IMG2D_ARRAY_RW_TY:]] = OpTypeImage %[[#VOID_TY]] 2D 0 1 0 0 Unknown ReadWrite
+; CHECK-SPIRV: %[[#IMG1D_WO_TY:]] = OpTypeImage %[[#VOID_TY]] 1D 0 0 0 0 Unknown WriteOnly
+; CHECK-SPIRV: %[[#IMG1D_RW_TY:]] = OpTypeImage %[[#VOID_TY]] 1D 0 0 0 0 Unknown ReadWrite
+; CHECK-SPIRV: %[[#IMG1D_BUFFER_WO_TY:]] = OpTypeImage %[[#VOID_TY]] Buffer 0 0 0 0 Unknown WriteOnly
+; CHECK-SPIRV: %[[#IMG1D_BUFFER_RW_TY:]] = OpTypeImage %[[#VOID_TY]] Buffer 0 0 0 0 Unknown ReadWrite
+; CHECK-SPIRV: %[[#IMG1D_ARRAY_WO_TY:]] = OpTypeImage %[[#VOID_TY]] 1D 0 1 0 0 Unknown WriteOnly
+; CHECK-SPIRV: %[[#IMG1D_ARRAY_RW_TY:]] = OpTypeImage %[[#VOID_TY]] 1D 0 1 0 0 Unknown ReadWrite
+; CHECK-SPIRV: %[[#IMG2D_DEPTH_WO_TY:]] = OpTypeImage %[[#VOID_TY]] 2D 1 0 0 0 Unknown WriteOnly
+; CHECK-SPIRV: %[[#IMG2D_ARRAY_DEPTH_WO_TY:]] = OpTypeImage %[[#VOID_TY]] 2D 1 1 0 0 Unknown WriteOnly
+; CHECK-SPIRV: %[[#IMG3D_WO_TY:]] = OpTypeImage %[[#VOID_TY]] 3D 0 0 0 0 Unknown WriteOnly
+; CHECK-SPIRV: %[[#IMG3D_RW_TY:]] = OpTypeImage %[[#VOID_TY]] 3D 0 0 0 0 Unknown ReadWrite
+
+%opencl.image2d_wo_t = type opaque
+%opencl.image2d_rw_t = type opaque
+%opencl.image2d_array_wo_t = type opaque
+%opencl.image2d_array_rw_t = type opaque
+%opencl.image1d_wo_t = type opaque
+%opencl.image1d_rw_t = type opaque
+%opencl.image1d_buffer_wo_t = type opaque
+%opencl.image1d_buffer_rw_t = type opaque
+%opencl.image1d_array_wo_t = type opaque
+%opencl.image1d_array_rw_t = type opaque
+%opencl.image2d_depth_wo_t = type opaque
+%opencl.image2d_array_depth_wo_t = type opaque
+%opencl.image3d_wo_t = type opaque
+%opencl.image3d_rw_t = type opaque
+
+;; kernel void test_img2d(write_only image2d_t image_wo, read_write image2d_t image_rw)
+;; {
+;;     write_imagef(image_wo, (int2)(0,0), (float4)(0,0,0,0));
+;;     write_imagei(image_wo, (int2)(0,0), (int4)(0,0,0,0));
+;;     write_imagef(image_rw, (int2)(0,0), (float4)(0,0,0,0));
+;;     write_imagei(image_rw, (int2)(0,0), (int4)(0,0,0,0));
+;;
+    ;; LOD
+;;     write_imagef(image_wo, (int2)(0,0), 0, (float4)(0,0,0,0));
+;;     write_imagei(image_wo, (int2)(0,0), 0, (int4)(0,0,0,0));
+;; }
+
+; CHECK-SPIRV: %[[#IMG2D_WO:]] = OpFunctionParameter %[[#IMG2D_WO_TY]]
+; CHECK-SPIRV: %[[#IMG2D_RW:]] = OpFunctionParameter %[[#IMG2D_RW_TY]]
+
+; CHECK-SPIRV: OpImageWrite %[[#IMG2D_WO]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG2D_WO]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG2D_RW]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG2D_RW]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG2D_WO]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG2D_WO]]
+
+define dso_local spir_kernel void @test_img2d(%opencl.image2d_wo_t addrspace(1)* %image_wo, %opencl.image2d_rw_t addrspace(1)* %image_rw) local_unnamed_addr {
+entry:
+  call spir_func void @_Z12write_imagef14ocl_image2d_woDv2_iDv4_f(%opencl.image2d_wo_t addrspace(1)* %image_wo, <2 x i32> noundef zeroinitializer, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei14ocl_image2d_woDv2_iDv4_i(%opencl.image2d_wo_t addrspace(1)* %image_wo, <2 x i32> noundef zeroinitializer, <4 x i32> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagef14ocl_image2d_rwDv2_iDv4_f(%opencl.image2d_rw_t addrspace(1)* %image_rw, <2 x i32> noundef zeroinitializer, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei14ocl_image2d_rwDv2_iDv4_i(%opencl.image2d_rw_t addrspace(1)* %image_rw, <2 x i32> noundef zeroinitializer, <4 x i32> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagef14ocl_image2d_woDv2_iiDv4_f(%opencl.image2d_wo_t addrspace(1)* %image_wo, <2 x i32> noundef zeroinitializer, i32 noundef 0, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei14ocl_image2d_woDv2_iiDv4_i(%opencl.image2d_wo_t addrspace(1)* %image_wo, <2 x i32> noundef zeroinitializer, i32 noundef 0, <4 x i32> noundef zeroinitializer)
+  ret void
+}
+
+declare spir_func void @_Z12write_imagef14ocl_image2d_woDv2_iDv4_f(%opencl.image2d_wo_t addrspace(1)*, <2 x i32> noundef, <4 x float> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagei14ocl_image2d_woDv2_iDv4_i(%opencl.image2d_wo_t addrspace(1)*, <2 x i32> noundef, <4 x i32> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagef14ocl_image2d_rwDv2_iDv4_f(%opencl.image2d_rw_t addrspace(1)*, <2 x i32> noundef, <4 x float> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagei14ocl_image2d_rwDv2_iDv4_i(%opencl.image2d_rw_t addrspace(1)*, <2 x i32> noundef, <4 x i32> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagef14ocl_image2d_woDv2_iiDv4_f(%opencl.image2d_wo_t addrspace(1)*, <2 x i32> noundef, i32 noundef, <4 x float> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagei14ocl_image2d_woDv2_iiDv4_i(%opencl.image2d_wo_t addrspace(1)*, <2 x i32> noundef, i32 noundef, <4 x i32> noundef) local_unnamed_addr
+
+;; kernel void test_img2d_array(write_only image2d_array_t image_wo, read_write image2d_array_t image_rw)
+;; {
+;;     write_imagef(image_wo, (int4)(0,0,0,0), (float4)(0,0,0,0));
+;;     write_imagei(image_wo, (int4)(0,0,0,0), (int4)(0,0,0,0));
+;;     write_imagef(image_rw, (int4)(0,0,0,0), (float4)(0,0,0,0));
+;;     write_imagei(image_rw, (int4)(0,0,0,0), (int4)(0,0,0,0));
+;;
+    ;; LOD
+;;     write_imagef(image_wo, (int4)(0,0,0,0), 0, (float4)(0,0,0,0));
+;;     write_imagei(image_wo, (int4)(0,0,0,0), 0, (int4)(0,0,0,0));
+;; }
+
+; CHECK-SPIRV: %[[#IMG2D_ARRAY_WO:]] = OpFunctionParameter %[[#IMG2D_ARRAY_WO_TY]]
+; CHECK-SPIRV: %[[#IMG2D_ARRAY_RW:]] = OpFunctionParameter %[[#IMG2D_ARRAY_RW_TY]]
+
+; CHECK-SPIRV: OpImageWrite %[[#IMG2D_ARRAY_WO]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG2D_ARRAY_WO]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG2D_ARRAY_RW]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG2D_ARRAY_RW]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG2D_ARRAY_WO]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG2D_ARRAY_WO]]
+
+define dso_local spir_kernel void @test_img2d_array(%opencl.image2d_array_wo_t addrspace(1)* %image_wo, %opencl.image2d_array_rw_t addrspace(1)* %image_rw) local_unnamed_addr {
+entry:
+  call spir_func void @_Z12write_imagef20ocl_image2d_array_woDv4_iDv4_f(%opencl.image2d_array_wo_t addrspace(1)* %image_wo, <4 x i32> noundef zeroinitializer, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei20ocl_image2d_array_woDv4_iS0_(%opencl.image2d_array_wo_t addrspace(1)* %image_wo, <4 x i32> noundef zeroinitializer, <4 x i32> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagef20ocl_image2d_array_rwDv4_iDv4_f(%opencl.image2d_array_rw_t addrspace(1)* %image_rw, <4 x i32> noundef zeroinitializer, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei20ocl_image2d_array_rwDv4_iS0_(%opencl.image2d_array_rw_t addrspace(1)* %image_rw, <4 x i32> noundef zeroinitializer, <4 x i32> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagef20ocl_image2d_array_woDv4_iiDv4_f(%opencl.image2d_array_wo_t addrspace(1)* %image_wo, <4 x i32> noundef zeroinitializer, i32 noundef 0, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei20ocl_image2d_array_woDv4_iiS0_(%opencl.image2d_array_wo_t addrspace(1)* %image_wo, <4 x i32> noundef zeroinitializer, i32 noundef 0, <4 x i32> noundef zeroinitializer)
+  ret void
+}
+
+declare spir_func void @_Z12write_imagef20ocl_image2d_array_woDv4_iDv4_f(%opencl.image2d_array_wo_t addrspace(1)*, <4 x i32> noundef, <4 x float> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagei20ocl_image2d_array_woDv4_iS0_(%opencl.image2d_array_wo_t addrspace(1)*, <4 x i32> noundef, <4 x i32> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagef20ocl_image2d_array_rwDv4_iDv4_f(%opencl.image2d_array_rw_t addrspace(1)*, <4 x i32> noundef, <4 x float> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagei20ocl_image2d_array_rwDv4_iS0_(%opencl.image2d_array_rw_t addrspace(1)*, <4 x i32> noundef, <4 x i32> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagef20ocl_image2d_array_woDv4_iiDv4_f(%opencl.image2d_array_wo_t addrspace(1)*, <4 x i32> noundef, i32 noundef, <4 x float> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagei20ocl_image2d_array_woDv4_iiS0_(%opencl.image2d_array_wo_t addrspace(1)*, <4 x i32> noundef, i32 noundef, <4 x i32> noundef) local_unnamed_addr
+
+;; kernel void test_img1d(write_only image1d_t image_wo, read_write image1d_t image_rw)
+;; {
+;;     write_imagef(image_wo, 0, (float4)(0,0,0,0));
+;;     write_imagei(image_wo, 0, (int4)(0,0,0,0));
+;;     write_imagef(image_rw, 0, (float4)(0,0,0,0));
+;;     write_imagei(image_rw, 0, (int4)(0,0,0,0));
+;;
+    ;; LOD
+;;     write_imagef(image_wo, 0, 0, (float4)(0,0,0,0));
+;;     write_imagei(image_wo, 0, 0, (int4)(0,0,0,0));
+;; }
+
+; CHECK-SPIRV: %[[#IMG1D_WO:]] = OpFunctionParameter %[[#IMG1D_WO_TY]]
+; CHECK-SPIRV: %[[#IMG1D_RW:]] = OpFunctionParameter %[[#IMG1D_RW_TY]]
+
+; CHECK-SPIRV: OpImageWrite %[[#IMG1D_WO]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG1D_WO]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG1D_RW]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG1D_RW]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG1D_WO]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG1D_WO]]
+
+define dso_local spir_kernel void @test_img1d(%opencl.image1d_wo_t addrspace(1)* %image_wo, %opencl.image1d_rw_t addrspace(1)* %image_rw) local_unnamed_addr {
+entry:
+  call spir_func void @_Z12write_imagef14ocl_image1d_woiDv4_f(%opencl.image1d_wo_t addrspace(1)* %image_wo, i32 noundef 0, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei14ocl_image1d_woiDv4_i(%opencl.image1d_wo_t addrspace(1)* %image_wo, i32 noundef 0, <4 x i32> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagef14ocl_image1d_rwiDv4_f(%opencl.image1d_rw_t addrspace(1)* %image_rw, i32 noundef 0, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei14ocl_image1d_rwiDv4_i(%opencl.image1d_rw_t addrspace(1)* %image_rw, i32 noundef 0, <4 x i32> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagef14ocl_image1d_woiiDv4_f(%opencl.image1d_wo_t addrspace(1)* %image_wo, i32 noundef 0, i32 noundef 0, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei14ocl_image1d_woiiDv4_i(%opencl.image1d_wo_t addrspace(1)* %image_wo, i32 noundef 0, i32 noundef 0, <4 x i32> noundef zeroinitializer)
+  ret void
+}
+
+declare spir_func void @_Z12write_imagef14ocl_image1d_woiDv4_f(%opencl.image1d_wo_t addrspace(1)*, i32 noundef, <4 x float> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagei14ocl_image1d_woiDv4_i(%opencl.image1d_wo_t addrspace(1)*, i32 noundef, <4 x i32> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagef14ocl_image1d_rwiDv4_f(%opencl.image1d_rw_t addrspace(1)*, i32 noundef, <4 x float> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagei14ocl_image1d_rwiDv4_i(%opencl.image1d_rw_t addrspace(1)*, i32 noundef, <4 x i32> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagef14ocl_image1d_woiiDv4_f(%opencl.image1d_wo_t addrspace(1)*, i32 noundef, i32 noundef, <4 x float> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagei14ocl_image1d_woiiDv4_i(%opencl.image1d_wo_t addrspace(1)*, i32 noundef, i32 noundef, <4 x i32> noundef) local_unnamed_addr
+
+;; kernel void test_img1d_buffer(write_only image1d_buffer_t image_wo, read_write image1d_buffer_t image_rw)
+;; {
+;;     write_imagef(image_wo, 0, (float4)(0,0,0,0));
+;;     write_imagei(image_wo, 0, (int4)(0,0,0,0));
+;;     write_imagef(image_rw, 0, (float4)(0,0,0,0));
+;;     write_imagei(image_rw, 0, (int4)(0,0,0,0));
+;; }
+
+; CHECK-SPIRV: %[[#IMG1D_BUFFER_WO:]] = OpFunctionParameter %[[#IMG1D_BUFFER_WO_TY]]
+; CHECK-SPIRV: %[[#IMG1D_BUFFER_RW:]] = OpFunctionParameter %[[#IMG1D_BUFFER_RW_TY]]
+
+; CHECK-SPIRV: OpImageWrite %[[#IMG1D_BUFFER_WO]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG1D_BUFFER_WO]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG1D_BUFFER_RW]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG1D_BUFFER_RW]]
+
+define dso_local spir_kernel void @test_img1d_buffer(%opencl.image1d_buffer_wo_t addrspace(1)* %image_wo, %opencl.image1d_buffer_rw_t addrspace(1)* %image_rw) local_unnamed_addr {
+entry:
+  call spir_func void @_Z12write_imagef21ocl_image1d_buffer_woiDv4_f(%opencl.image1d_buffer_wo_t addrspace(1)* %image_wo, i32 noundef 0, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei21ocl_image1d_buffer_woiDv4_i(%opencl.image1d_buffer_wo_t addrspace(1)* %image_wo, i32 noundef 0, <4 x i32> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagef21ocl_image1d_buffer_rwiDv4_f(%opencl.image1d_buffer_rw_t addrspace(1)* %image_rw, i32 noundef 0, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei21ocl_image1d_buffer_rwiDv4_i(%opencl.image1d_buffer_rw_t addrspace(1)* %image_rw, i32 noundef 0, <4 x i32> noundef zeroinitializer)
+  ret void
+}
+
+declare spir_func void @_Z12write_imagef21ocl_image1d_buffer_woiDv4_f(%opencl.image1d_buffer_wo_t addrspace(1)*, i32 noundef, <4 x float> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagei21ocl_image1d_buffer_woiDv4_i(%opencl.image1d_buffer_wo_t addrspace(1)*, i32 noundef, <4 x i32> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagef21ocl_image1d_buffer_rwiDv4_f(%opencl.image1d_buffer_rw_t addrspace(1)*, i32 noundef, <4 x float> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagei21ocl_image1d_buffer_rwiDv4_i(%opencl.image1d_buffer_rw_t addrspace(1)*, i32 noundef, <4 x i32> noundef) local_unnamed_addr
+
+;; kernel void test_img1d_array(write_only image1d_array_t image_wo, read_write image1d_array_t image_rw)
+;; {
+;;     write_imagef(image_wo, (int2)(0,0), (float4)(0,0,0,0));
+;;     write_imagei(image_wo, (int2)(0,0), (int4)(0,0,0,0));
+;;     write_imagef(image_rw, (int2)(0,0), (float4)(0,0,0,0));
+;;     write_imagei(image_rw, (int2)(0,0), (int4)(0,0,0,0));
+;;
+    ;; LOD
+;;     write_imagef(image_wo, (int2)(0,0), 0, (float4)(0,0,0,0));
+;;     write_imagei(image_wo, (int2)(0,0), 0, (int4)(0,0,0,0));
+;; }
+
+; CHECK-SPIRV: %[[#IMG1D_ARRAY_WO:]] = OpFunctionParameter %[[#IMG1D_ARRAY_WO_TY]]
+; CHECK-SPIRV: %[[#IMG1D_ARRAY_RW:]] = OpFunctionParameter %[[#IMG1D_ARRAY_RW_TY]]
+
+; CHECK-SPIRV: OpImageWrite %[[#IMG1D_ARRAY_WO]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG1D_ARRAY_WO]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG1D_ARRAY_RW]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG1D_ARRAY_RW]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG1D_ARRAY_WO]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG1D_ARRAY_WO]]
+
+define dso_local spir_kernel void @test_img1d_array(%opencl.image1d_array_wo_t addrspace(1)* %image_wo, %opencl.image1d_array_rw_t addrspace(1)* %image_rw) local_unnamed_addr {
+entry:
+  call spir_func void @_Z12write_imagef20ocl_image1d_array_woDv2_iDv4_f(%opencl.image1d_array_wo_t addrspace(1)* %image_wo, <2 x i32> noundef zeroinitializer, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei20ocl_image1d_array_woDv2_iDv4_i(%opencl.image1d_array_wo_t addrspace(1)* %image_wo, <2 x i32> noundef zeroinitializer, <4 x i32> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagef20ocl_image1d_array_rwDv2_iDv4_f(%opencl.image1d_array_rw_t addrspace(1)* %image_rw, <2 x i32> noundef zeroinitializer, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei20ocl_image1d_array_rwDv2_iDv4_i(%opencl.image1d_array_rw_t addrspace(1)* %image_rw, <2 x i32> noundef zeroinitializer, <4 x i32> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagef20ocl_image1d_array_woDv2_iiDv4_f(%opencl.image1d_array_wo_t addrspace(1)* %image_wo, <2 x i32> noundef zeroinitializer, i32 noundef 0, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei20ocl_image1d_array_woDv2_iiDv4_i(%opencl.image1d_array_wo_t addrspace(1)* %image_wo, <2 x i32> noundef zeroinitializer, i32 noundef 0, <4 x i32> noundef zeroinitializer)
+  ret void
+}
+
+declare spir_func void @_Z12write_imagef20ocl_image1d_array_woDv2_iDv4_f(%opencl.image1d_array_wo_t addrspace(1)*, <2 x i32> noundef, <4 x float> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagei20ocl_image1d_array_woDv2_iDv4_i(%opencl.image1d_array_wo_t addrspace(1)*, <2 x i32> noundef, <4 x i32> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagef20ocl_image1d_array_rwDv2_iDv4_f(%opencl.image1d_array_rw_t addrspace(1)*, <2 x i32> noundef, <4 x float> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagei20ocl_image1d_array_rwDv2_iDv4_i(%opencl.image1d_array_rw_t addrspace(1)*, <2 x i32> noundef, <4 x i32> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagef20ocl_image1d_array_woDv2_iiDv4_f(%opencl.image1d_array_wo_t addrspace(1)*, <2 x i32> noundef, i32 noundef, <4 x float> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagei20ocl_image1d_array_woDv2_iiDv4_i(%opencl.image1d_array_wo_t addrspace(1)*, <2 x i32> noundef, i32 noundef, <4 x i32> noundef) local_unnamed_addr
+
+;; kernel void test_img2d_depth(write_only image2d_depth_t image_wo)
+;; {
+;;     write_imagef(image_wo, (int2)(0,0), (float)(0));
+;;     write_imagef(image_wo, (int2)(0,0), (float)(0));
+;;
+    ;; LOD
+;;     write_imagef(image_wo, (int2)(0,0), 0, (float)(0));
+;; }
+
+; CHECK-SPIRV: %[[#IMG2D_DEPTH_WO:]] = OpFunctionParameter %[[#IMG2D_DEPTH_WO_TY]]
+
+; CHECK-SPIRV: OpImageWrite %[[#IMG2D_DEPTH_WO]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG2D_DEPTH_WO]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG2D_DEPTH_WO]]
+
+define dso_local spir_kernel void @test_img2d_depth(%opencl.image2d_depth_wo_t addrspace(1)* %image_wo) local_unnamed_addr {
+entry:
+  call spir_func void @_Z12write_imagef20ocl_image2d_depth_woDv2_if(%opencl.image2d_depth_wo_t addrspace(1)* %image_wo, <2 x i32> noundef zeroinitializer, float noundef 0.000000e+00)
+  call spir_func void @_Z12write_imagef20ocl_image2d_depth_woDv2_if(%opencl.image2d_depth_wo_t addrspace(1)* %image_wo, <2 x i32> noundef zeroinitializer, float noundef 0.000000e+00)
+  call spir_func void @_Z12write_imagef20ocl_image2d_depth_woDv2_iif(%opencl.image2d_depth_wo_t addrspace(1)* %image_wo, <2 x i32> noundef zeroinitializer, i32 noundef 0, float noundef 0.000000e+00)
+  ret void
+}
+
+declare spir_func void @_Z12write_imagef20ocl_image2d_depth_woDv2_if(%opencl.image2d_depth_wo_t addrspace(1)*, <2 x i32> noundef, float noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagef20ocl_image2d_depth_woDv2_iif(%opencl.image2d_depth_wo_t addrspace(1)*, <2 x i32> noundef, i32 noundef, float noundef) local_unnamed_addr
+
+;; kernel void test_img2d_array_depth(write_only image2d_array_depth_t image_wo)
+;; {
+;;     write_imagef(image_wo, (int4)(0,0,0,0), (float)(0));
+;;     write_imagef(image_wo, (int4)(0,0,0,0), (float)(0));
+;;
+    ;; LOD
+;;     write_imagef(image_wo, (int4)(0,0,0,0), 0, (float)(0));
+;; }
+
+; CHECK-SPIRV: %[[#IMG2D_ARRAY_DEPTH_WO:]] = OpFunctionParameter %[[#IMG2D_ARRAY_DEPTH_WO_TY]]
+
+; CHECK-SPIRV: OpImageWrite %[[#IMG2D_ARRAY_DEPTH_WO]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG2D_ARRAY_DEPTH_WO]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG2D_ARRAY_DEPTH_WO]]
+
+define dso_local spir_kernel void @test_img2d_array_depth(%opencl.image2d_array_depth_wo_t addrspace(1)* %image_wo) local_unnamed_addr {
+entry:
+  call spir_func void @_Z12write_imagef26ocl_image2d_array_depth_woDv4_if(%opencl.image2d_array_depth_wo_t addrspace(1)* %image_wo, <4 x i32> noundef zeroinitializer, float noundef 0.000000e+00)
+  call spir_func void @_Z12write_imagef26ocl_image2d_array_depth_woDv4_if(%opencl.image2d_array_depth_wo_t addrspace(1)* %image_wo, <4 x i32> noundef zeroinitializer, float noundef 0.000000e+00)
+  call spir_func void @_Z12write_imagef26ocl_image2d_array_depth_woDv4_iif(%opencl.image2d_array_depth_wo_t addrspace(1)* %image_wo, <4 x i32> noundef zeroinitializer, i32 noundef 0, float noundef 0.000000e+00)
+  ret void
+}
+
+declare spir_func void @_Z12write_imagef26ocl_image2d_array_depth_woDv4_if(%opencl.image2d_array_depth_wo_t addrspace(1)*, <4 x i32> noundef, float noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagef26ocl_image2d_array_depth_woDv4_iif(%opencl.image2d_array_depth_wo_t addrspace(1)*, <4 x i32> noundef, i32 noundef, float noundef) local_unnamed_addr
+
+;; kernel void test_img3d(write_only image3d_t image_wo, read_write image3d_t image_rw)
+;; {
+;;     write_imagef(image_wo, (int4)(0,0,0,0), (float4)(0,0,0,0));
+;;     write_imagei(image_wo, (int4)(0,0,0,0), (int4)(0,0,0,0));
+;;     write_imagef(image_rw, (int4)(0,0,0,0), (float4)(0,0,0,0));
+;;     write_imagei(image_rw, (int4)(0,0,0,0), (int4)(0,0,0,0));
+;;
+    ;; LOD
+;;     write_imagef(image_wo, (int4)(0,0,0,0), 0, (float4)(0,0,0,0));
+;;     write_imagei(image_wo, (int4)(0,0,0,0), 0, (int4)(0,0,0,0));
+;; }
+
+; CHECK-SPIRV: %[[#IMG3D_WO:]] = OpFunctionParameter %[[#IMG3D_WO_TY]]
+; CHECK-SPIRV: %[[#IMG3D_RW:]] = OpFunctionParameter %[[#IMG3D_RW_TY]]
+
+; CHECK-SPIRV: OpImageWrite %[[#IMG3D_WO]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG3D_WO]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG3D_RW]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG3D_RW]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG3D_WO]]
+; CHECK-SPIRV: OpImageWrite %[[#IMG3D_WO]]
+
+define dso_local spir_kernel void @test_img3d(%opencl.image3d_wo_t addrspace(1)* %image_wo, %opencl.image3d_rw_t addrspace(1)* %image_rw) local_unnamed_addr {
+entry:
+  call spir_func void @_Z12write_imagef14ocl_image3d_woDv4_iDv4_f(%opencl.image3d_wo_t addrspace(1)* %image_wo, <4 x i32> noundef zeroinitializer, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei14ocl_image3d_woDv4_iS0_(%opencl.image3d_wo_t addrspace(1)* %image_wo, <4 x i32> noundef zeroinitializer, <4 x i32> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagef14ocl_image3d_rwDv4_iDv4_f(%opencl.image3d_rw_t addrspace(1)* %image_rw, <4 x i32> noundef zeroinitializer, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei14ocl_image3d_rwDv4_iS0_(%opencl.image3d_rw_t addrspace(1)* %image_rw, <4 x i32> noundef zeroinitializer, <4 x i32> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagef14ocl_image3d_woDv4_iiDv4_f(%opencl.image3d_wo_t addrspace(1)* %image_wo, <4 x i32> noundef zeroinitializer, i32 noundef 0, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei14ocl_image3d_woDv4_iiS0_(%opencl.image3d_wo_t addrspace(1)* %image_wo, <4 x i32> noundef zeroinitializer, i32 noundef 0, <4 x i32> noundef zeroinitializer)
+  ret void
+}
+
+declare spir_func void @_Z12write_imagef14ocl_image3d_woDv4_iDv4_f(%opencl.image3d_wo_t addrspace(1)*, <4 x i32> noundef, <4 x float> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagei14ocl_image3d_woDv4_iS0_(%opencl.image3d_wo_t addrspace(1)*, <4 x i32> noundef, <4 x i32> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagef14ocl_image3d_rwDv4_iDv4_f(%opencl.image3d_rw_t addrspace(1)*, <4 x i32> noundef, <4 x float> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagei14ocl_image3d_rwDv4_iS0_(%opencl.image3d_rw_t addrspace(1)*, <4 x i32> noundef, <4 x i32> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagef14ocl_image3d_woDv4_iiDv4_f(%opencl.image3d_wo_t addrspace(1)*, <4 x i32> noundef, i32 noundef, <4 x float> noundef) local_unnamed_addr
+
+declare spir_func void @_Z12write_imagei14ocl_image3d_woDv4_iiS0_(%opencl.image3d_wo_t addrspace(1)*, <4 x i32> noundef, i32 noundef, <4 x i32> noundef) local_unnamed_addr

diff  --git a/llvm/test/CodeGen/SPIRV/transcoding/OpSwitchUnreachable.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpSwitchUnreachable.ll
new file mode 100644
index 0000000000000..e73efbeade70d
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpSwitchUnreachable.ll
@@ -0,0 +1,20 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+define void @test_switch_with_unreachable_block(i1 %a) {
+  %value = zext i1 %a to i32
+; CHECK-SPIRV:      OpSwitch %[[#]] %[[#REACHABLE:]]
+  switch i32 %value, label %unreachable [
+    i32 0, label %reachable
+    i32 1, label %reachable
+  ]
+
+; CHECK-SPIRV-NEXT: %[[#REACHABLE]] = OpLabel
+reachable:
+; CHECK-SPIRV-NEXT: OpReturn
+  ret void
+
+; CHECK-SPIRV:      %[[#]] = OpLabel
+; CHECK-SPIRV-NEXT: OpUnreachable
+unreachable:
+  unreachable
+}

diff  --git a/llvm/test/CodeGen/SPIRV/transcoding/SampledImage.ll b/llvm/test/CodeGen/SPIRV/transcoding/SampledImage.ll
new file mode 100644
index 0000000000000..73981becc3b3f
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/SampledImage.ll
@@ -0,0 +1,89 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+;; constant sampler_t constSampl = CLK_FILTER_LINEAR;
+;;
+;; __kernel
+;; void sample_kernel_float(image2d_t input, float2 coords, global float4 *results, sampler_t argSampl) {
+;;   *results = read_imagef(input, constSampl, coords);
+;;   *results = read_imagef(input, argSampl, coords);
+;;   *results = read_imagef(input, CLK_FILTER_NEAREST|CLK_ADDRESS_REPEAT, coords);
+;; }
+;;
+;; __kernel
+;; void sample_kernel_int(image2d_t input, float2 coords, global int4 *results, sampler_t argSampl) {
+;;   *results = read_imagei(input, constSampl, coords);
+;;   *results = read_imagei(input, argSampl, coords);
+;;   *results = read_imagei(input, CLK_FILTER_NEAREST|CLK_ADDRESS_REPEAT, coords);
+;; }
+
+%opencl.image2d_ro_t = type opaque
+%opencl.sampler_t = type opaque
+
+; CHECK-SPIRV: OpCapability LiteralSampler
+; CHECK-SPIRV: OpName %[[#sample_kernel_float:]] "sample_kernel_float"
+; CHECK-SPIRV: OpName %[[#sample_kernel_int:]] "sample_kernel_int"
+
+; CHECK-SPIRV:     %[[#TypeSampler:]] = OpTypeSampler
+; CHECK-SPIRV-DAG: %[[#SampledImageTy:]] = OpTypeSampledImage
+; CHECK-SPIRV-DAG: %[[#ConstSampler1:]] = OpConstantSampler %[[#TypeSampler]] None 0 Linear
+; CHECK-SPIRV-DAG: %[[#ConstSampler2:]] = OpConstantSampler %[[#TypeSampler]] Repeat 0 Nearest
+; CHECK-SPIRV-DAG: %[[#ConstSampler3:]] = OpConstantSampler %[[#TypeSampler]] None 0 Linear
+; CHECK-SPIRV-DAG: %[[#ConstSampler4:]] = OpConstantSampler %[[#TypeSampler]] Repeat 0 Nearest
+
+; CHECK-SPIRV: %[[#sample_kernel_float]] = OpFunction %{{.*}}
+; CHECK-SPIRV: %[[#InputImage:]] = OpFunctionParameter %{{.*}}
+; CHECK-SPIRV: %[[#argSampl:]] = OpFunctionParameter %[[#TypeSampler]]
+
+; CHECK-SPIRV: %[[#SampledImage1:]] = OpSampledImage %[[#SampledImageTy]] %[[#InputImage]] %[[#ConstSampler1]]
+; CHECK-SPIRV: %[[#]] = OpImageSampleExplicitLod %[[#]] %[[#SampledImage1]]
+
+; CHECK-SPIRV: %[[#SampledImage2:]] = OpSampledImage %[[#SampledImageTy]] %[[#InputImage]] %[[#argSampl]]
+; CHECK-SPIRV: %[[#]] = OpImageSampleExplicitLod %[[#]] %[[#SampledImage2]]
+
+; CHECK-SPIRV: %[[#SampledImage3:]] = OpSampledImage %[[#SampledImageTy]] %[[#InputImage]] %[[#ConstSampler2]]
+; CHECK-SPIRV: %[[#]] = OpImageSampleExplicitLod %[[#]] %[[#SampledImage3]]
+
+define dso_local spir_kernel void @sample_kernel_float(%opencl.image2d_ro_t addrspace(1)* %input, <2 x float> noundef %coords, <4 x float> addrspace(1)* nocapture noundef writeonly %results, %opencl.sampler_t addrspace(2)* %argSampl) local_unnamed_addr {
+entry:
+  %0 = tail call spir_func %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 32)
+  %call = tail call spir_func <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv2_f(%opencl.image2d_ro_t addrspace(1)* %input, %opencl.sampler_t addrspace(2)* %0, <2 x float> noundef %coords)
+  store <4 x float> %call, <4 x float> addrspace(1)* %results, align 16
+  %call1 = tail call spir_func <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv2_f(%opencl.image2d_ro_t addrspace(1)* %input, %opencl.sampler_t addrspace(2)* %argSampl, <2 x float> noundef %coords)
+  store <4 x float> %call1, <4 x float> addrspace(1)* %results, align 16
+  %1 = tail call spir_func %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 22)
+  %call2 = tail call spir_func <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv2_f(%opencl.image2d_ro_t addrspace(1)* %input, %opencl.sampler_t addrspace(2)* %1, <2 x float> noundef %coords)
+  store <4 x float> %call2, <4 x float> addrspace(1)* %results, align 16
+  ret void
+}
+
+declare spir_func <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv2_f(%opencl.image2d_ro_t addrspace(1)*, %opencl.sampler_t addrspace(2)*, <2 x float> noundef) local_unnamed_addr
+
+declare spir_func %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32) local_unnamed_addr
+
+; CHECK-SPIRV: %[[#sample_kernel_int]] = OpFunction %{{.*}}
+; CHECK-SPIRV: %[[#InputImage:]] = OpFunctionParameter %{{.*}}
+; CHECK-SPIRV: %[[#argSampl:]] = OpFunctionParameter %[[#TypeSampler]]
+
+; CHECK-SPIRV: %[[#SampledImage4:]] = OpSampledImage %[[#SampledImageTy]] %[[#InputImage]] %[[#ConstSampler3]]
+; CHECK-SPIRV: %[[#]] = OpImageSampleExplicitLod %[[#]] %[[#SampledImage4]]
+
+; CHECK-SPIRV: %[[#SampledImage5:]] = OpSampledImage %[[#SampledImageTy]] %[[#InputImage]] %[[#argSampl]]
+; CHECK-SPIRV: %[[#]] = OpImageSampleExplicitLod %[[#]] %[[#SampledImage5]]
+
+; CHECK-SPIRV: %[[#SampledImage6:]] = OpSampledImage %[[#SampledImageTy]] %[[#InputImage]] %[[#ConstSampler4]]
+; CHECK-SPIRV: %[[#]] = OpImageSampleExplicitLod %[[#]] %[[#SampledImage6]]
+
+define dso_local spir_kernel void @sample_kernel_int(%opencl.image2d_ro_t addrspace(1)* %input, <2 x float> noundef %coords, <4 x i32> addrspace(1)* nocapture noundef writeonly %results, %opencl.sampler_t addrspace(2)* %argSampl) local_unnamed_addr {
+entry:
+  %0 = tail call spir_func %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 32)
+  %call = tail call spir_func <4 x i32> @_Z11read_imagei14ocl_image2d_ro11ocl_samplerDv2_f(%opencl.image2d_ro_t addrspace(1)* %input, %opencl.sampler_t addrspace(2)* %0, <2 x float> noundef %coords)
+  store <4 x i32> %call, <4 x i32> addrspace(1)* %results, align 16
+  %call1 = tail call spir_func <4 x i32> @_Z11read_imagei14ocl_image2d_ro11ocl_samplerDv2_f(%opencl.image2d_ro_t addrspace(1)* %input, %opencl.sampler_t addrspace(2)* %argSampl, <2 x float> noundef %coords)
+  store <4 x i32> %call1, <4 x i32> addrspace(1)* %results, align 16
+  %1 = tail call spir_func %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 22)
+  %call2 = tail call spir_func <4 x i32> @_Z11read_imagei14ocl_image2d_ro11ocl_samplerDv2_f(%opencl.image2d_ro_t addrspace(1)* %input, %opencl.sampler_t addrspace(2)* %1, <2 x float> noundef %coords)
+  store <4 x i32> %call2, <4 x i32> addrspace(1)* %results, align 16
+  ret void
+}
+
+declare spir_func <4 x i32> @_Z11read_imagei14ocl_image2d_ro11ocl_samplerDv2_f(%opencl.image2d_ro_t addrspace(1)*, %opencl.sampler_t addrspace(2)*, <2 x float> noundef) local_unnamed_addr

diff  --git a/llvm/test/CodeGen/SPIRV/transcoding/check_ro_qualifier.ll b/llvm/test/CodeGen/SPIRV/transcoding/check_ro_qualifier.ll
new file mode 100644
index 0000000000000..731fe55b50774
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/check_ro_qualifier.ll
@@ -0,0 +1,23 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: %[[#IMAGE_TYPE:]] = OpTypeImage
+; CHECK-SPIRV: %[[#IMAGE_ARG:]] = OpFunctionParameter %[[#IMAGE_TYPE]]
+; CHECK-SPIRV: %[[#]] = OpImageQuerySizeLod %[[#]] %[[#IMAGE_ARG]]
+; CHECK-SPIRV: %[[#]] = OpImageQuerySizeLod %[[#]] %[[#IMAGE_ARG]]
+
+%opencl.image2d_array_ro_t = type opaque
+
+define spir_kernel void @sample_kernel(%opencl.image2d_array_ro_t addrspace(1)* %input) {
+entry:
+  %call.tmp1 = call spir_func <2 x i32> @_Z13get_image_dim20ocl_image2d_array_ro(%opencl.image2d_array_ro_t addrspace(1)* %input)
+  %call.tmp2 = shufflevector <2 x i32> %call.tmp1, <2 x i32> undef, <3 x i32> <i32 0, i32 1, i32 2>
+  %call.tmp3 = call spir_func i64 @_Z20get_image_array_size20ocl_image2d_array_ro(%opencl.image2d_array_ro_t addrspace(1)* %input)
+  %call.tmp4 = trunc i64 %call.tmp3 to i32
+  %call.tmp5 = insertelement <3 x i32> %call.tmp2, i32 %call.tmp4, i32 2
+  %call.old = extractelement <3 x i32> %call.tmp5, i32 0
+  ret void
+}
+
+declare spir_func <2 x i32> @_Z13get_image_dim20ocl_image2d_array_ro(%opencl.image2d_array_ro_t addrspace(1)*)
+
+declare spir_func i64 @_Z20get_image_array_size20ocl_image2d_array_ro(%opencl.image2d_array_ro_t addrspace(1)*)

diff  --git a/llvm/test/CodeGen/SPIRV/transcoding/cl-types.ll b/llvm/test/CodeGen/SPIRV/transcoding/cl-types.ll
new file mode 100644
index 0000000000000..19910ff1d11b0
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/cl-types.ll
@@ -0,0 +1,87 @@
+;; Test CL opaque types
+;;
+;; // cl-types.cl
+;; // CL source code for generating LLVM IR.
+;; // Command for compilation:
+;; //  clang -cc1 -x cl -cl-std=CL2.0 -triple spir-unknown-unknown -emit-llvm cl-types.cl
+;; void kernel foo(
+;;  read_only pipe int a,
+;;  write_only pipe int b,
+;;  read_only image1d_t c1,
+;;  read_only image2d_t d1,
+;;  read_only image3d_t e1,
+;;  read_only image2d_array_t f1,
+;;  read_only image1d_buffer_t g1,
+;;  write_only image1d_t c2,
+;;  read_write image2d_t d3,
+;;  sampler_t s
+;; ) {
+;; }
+
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV-DAG: OpCapability Sampled1D
+; CHECK-SPIRV-DAG: OpCapability SampledBuffer
+; CHECK-SPIRV-DAG: %[[#VOID:]] = OpTypeVoid
+; CHECK-SPIRV-DAG: %[[#PIPE_RD:]] = OpTypePipe ReadOnly
+; CHECK-SPIRV-DAG: %[[#PIPE_WR:]] = OpTypePipe WriteOnly
+; CHECK-SPIRV-DAG: %[[#IMG1D_RD:]] = OpTypeImage %[[#VOID]] 1D 0 0 0 0 Unknown ReadOnly
+; CHECK-SPIRV-DAG: %[[#IMG2D_RD:]] = OpTypeImage %[[#VOID]] 2D 0 0 0 0 Unknown ReadOnly
+; CHECK-SPIRV-DAG: %[[#IMG3D_RD:]] = OpTypeImage %[[#VOID]] 3D 0 0 0 0 Unknown ReadOnly
+; CHECK-SPIRV-DAG: %[[#IMG2DA_RD:]] = OpTypeImage %[[#VOID]] 2D 0 1 0 0 Unknown ReadOnly
+; CHECK-SPIRV-DAG: %[[#IMG1DB_RD:]] = OpTypeImage %[[#VOID]] Buffer 0 0 0 0 Unknown ReadOnly
+; CHECK-SPIRV-DAG: %[[#IMG1D_WR:]] = OpTypeImage %[[#VOID]] 1D 0 0 0 0 Unknown WriteOnly
+; CHECK-SPIRV-DAG: %[[#IMG2D_RW:]] = OpTypeImage %[[#VOID]] 2D 0 0 0 0 Unknown ReadWrite
+; CHECK-SPIRV-DAG: %[[#SAMP:]] = OpTypeSampler
+; CHECK-SPIRV-DAG: %[[#SAMPIMG:]] = OpTypeSampledImage %[[#IMG2D_RD]]
+
+; CHECK-SPIRV:     %[[#SAMP_CONST:]] = OpConstantSampler %[[#SAMP]] None 0 Linear
+
+%opencl.pipe_ro_t = type opaque
+%opencl.pipe_wo_t = type opaque
+%opencl.image3d_ro_t = type opaque
+%opencl.image2d_array_ro_t = type opaque
+%opencl.image1d_buffer_ro_t = type opaque
+%opencl.image1d_ro_t = type opaque
+%opencl.image1d_wo_t = type opaque
+%opencl.image2d_rw_t = type opaque
+%opencl.image2d_ro_t = type opaque
+%opencl.sampler_t = type opaque
+
+; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#PIPE_RD]]
+; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#PIPE_WR]]
+; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#IMG1D_RD]]
+; CHECK-SPIRV: %[[#IMG_ARG:]] = OpFunctionParameter %[[#IMG2D_RD]]
+; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#IMG3D_RD]]
+; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#IMG2DA_RD]]
+; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#IMG1DB_RD]]
+; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#IMG1D_WR]]
+; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#IMG2D_RW]]
+; CHECK-SPIRV: %[[#SAMP_ARG:]] = OpFunctionParameter %[[#SAMP]]
+
+define spir_kernel void @foo(
+  %opencl.pipe_ro_t addrspace(1)* nocapture %a,
+  %opencl.pipe_wo_t addrspace(1)* nocapture %b,
+  %opencl.image1d_ro_t addrspace(1)* nocapture %c1,
+  %opencl.image2d_ro_t addrspace(1)* nocapture %d1,
+  %opencl.image3d_ro_t addrspace(1)* nocapture %e1,
+  %opencl.image2d_array_ro_t addrspace(1)* nocapture %f1,
+  %opencl.image1d_buffer_ro_t addrspace(1)* nocapture %g1,
+  %opencl.image1d_wo_t addrspace(1)* nocapture %c2,
+  %opencl.image2d_rw_t addrspace(1)* nocapture %d3,
+  %opencl.sampler_t addrspace(2)* %s) {
+entry:
+; CHECK-SPIRV: %[[#SAMPIMG_VAR1:]] = OpSampledImage %[[#SAMPIMG]] %[[#IMG_ARG]] %[[#SAMP_ARG]]
+; CHECK-SPIRV: %[[#]] = OpImageSampleExplicitLod %[[#]] %[[#SAMPIMG_VAR1]]
+  %.tmp = call spir_func <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv4_if(%opencl.image2d_ro_t addrspace(1)* %d1, %opencl.sampler_t addrspace(2)* %s, <4 x i32> zeroinitializer, float 1.000000e+00)
+
+; CHECK-SPIRV: %[[#SAMPIMG_VAR2:]] = OpSampledImage %[[#SAMPIMG]] %[[#IMG_ARG]] %[[#SAMP_CONST]]
+; CHECK-SPIRV: %[[#]] = OpImageSampleExplicitLod %[[#]] %[[#SAMPIMG_VAR2]]
+  %0 = call %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 32)
+  %.tmp2 = call spir_func <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv4_if(%opencl.image2d_ro_t addrspace(1)* %d1, %opencl.sampler_t addrspace(2)* %0, <4 x i32> zeroinitializer, float 1.000000e+00)
+  ret void
+}
+
+declare spir_func <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv4_if(%opencl.image2d_ro_t addrspace(1)*, %opencl.sampler_t addrspace(2)*, <4 x i32>, float)
+
+declare %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32)

diff  --git a/llvm/test/CodeGen/SPIRV/transcoding/clk_event_t.ll b/llvm/test/CodeGen/SPIRV/transcoding/clk_event_t.ll
new file mode 100644
index 0000000000000..1f7e981a07cee
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/clk_event_t.ll
@@ -0,0 +1,47 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: OpTypeDeviceEvent
+; CHECK-SPIRV: OpFunction
+; CHECK-SPIRV: OpCreateUserEvent
+; CHECK-SPIRV: OpIsValidEvent
+; CHECK-SPIRV: OpRetainEvent
+; CHECK-SPIRV: OpSetUserEventStatus
+; CHECK-SPIRV: OpCaptureEventProfilingInfo
+; CHECK-SPIRV: OpReleaseEvent
+; CHECK-SPIRV: OpFunctionEnd
+
+;; kernel void clk_event_t_test(global int *res, global void *prof) {
+;;   clk_event_t e1 = create_user_event();
+;;   *res = is_valid_event(e1);
+;;   retain_event(e1);
+;;   set_user_event_status(e1, -42);
+;;   capture_event_profiling_info(e1, CLK_PROFILING_COMMAND_EXEC_TIME, prof);
+;;   release_event(e1);
+;; }
+
+%opencl.clk_event_t = type opaque
+
+define dso_local spir_kernel void @clk_event_t_test(i32 addrspace(1)* nocapture noundef writeonly %res, i8 addrspace(1)* noundef %prof) local_unnamed_addr {
+entry:
+  %call = call spir_func %opencl.clk_event_t* @_Z17create_user_eventv()
+  %call1 = call spir_func zeroext i1 @_Z14is_valid_event12ocl_clkevent(%opencl.clk_event_t* %call)
+  %conv = zext i1 %call1 to i32
+  store i32 %conv, i32 addrspace(1)* %res, align 4
+  call spir_func void @_Z12retain_event12ocl_clkevent(%opencl.clk_event_t* %call)
+  call spir_func void @_Z21set_user_event_status12ocl_clkeventi(%opencl.clk_event_t* %call, i32 noundef -42)
+  call spir_func void @_Z28capture_event_profiling_info12ocl_clkeventiPU3AS1v(%opencl.clk_event_t* %call, i32 noundef 1, i8 addrspace(1)* noundef %prof)
+  call spir_func void @_Z13release_event12ocl_clkevent(%opencl.clk_event_t* %call)
+  ret void
+}
+
+declare spir_func %opencl.clk_event_t* @_Z17create_user_eventv() local_unnamed_addr
+
+declare spir_func zeroext i1 @_Z14is_valid_event12ocl_clkevent(%opencl.clk_event_t*) local_unnamed_addr
+
+declare spir_func void @_Z12retain_event12ocl_clkevent(%opencl.clk_event_t*) local_unnamed_addr
+
+declare spir_func void @_Z21set_user_event_status12ocl_clkeventi(%opencl.clk_event_t*, i32 noundef) local_unnamed_addr
+
+declare spir_func void @_Z28capture_event_profiling_info12ocl_clkeventiPU3AS1v(%opencl.clk_event_t*, i32 noundef, i8 addrspace(1)* noundef) local_unnamed_addr
+
+declare spir_func void @_Z13release_event12ocl_clkevent(%opencl.clk_event_t*) local_unnamed_addr

diff  --git a/llvm/test/CodeGen/SPIRV/transcoding/explicit-conversions.ll b/llvm/test/CodeGen/SPIRV/transcoding/explicit-conversions.ll
new file mode 100644
index 0000000000000..c186a8135fee7
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/explicit-conversions.ll
@@ -0,0 +1,117 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: OpSatConvertSToU
+
+;; kernel void testSToU(global int2 *a, global uchar2 *res) {
+;;   res[0] = convert_uchar2_sat(*a);
+;; }
+
+define dso_local spir_kernel void @testSToU(<2 x i32> addrspace(1)* nocapture noundef readonly %a, <2 x i8> addrspace(1)* nocapture noundef writeonly %res) local_unnamed_addr {
+entry:
+  %0 = load <2 x i32>, <2 x i32> addrspace(1)* %a, align 8
+  %call = call spir_func <2 x i8> @_Z18convert_uchar2_satDv2_i(<2 x i32> noundef %0)
+  store <2 x i8> %call, <2 x i8> addrspace(1)* %res, align 2
+  ret void
+}
+
+declare spir_func <2 x i8> @_Z18convert_uchar2_satDv2_i(<2 x i32> noundef) local_unnamed_addr
+
+; CHECK-SPIRV: OpSatConvertUToS
+
+;; kernel void testUToS(global uint2 *a, global char2 *res) {
+;;   res[0] = convert_char2_sat(*a);
+;; }
+
+define dso_local spir_kernel void @testUToS(<2 x i32> addrspace(1)* nocapture noundef readonly %a, <2 x i8> addrspace(1)* nocapture noundef writeonly %res) local_unnamed_addr {
+entry:
+  %0 = load <2 x i32>, <2 x i32> addrspace(1)* %a, align 8
+  %call = call spir_func <2 x i8> @_Z17convert_char2_satDv2_j(<2 x i32> noundef %0)
+  store <2 x i8> %call, <2 x i8> addrspace(1)* %res, align 2
+  ret void
+}
+
+declare spir_func <2 x i8> @_Z17convert_char2_satDv2_j(<2 x i32> noundef) local_unnamed_addr
+
+; CHECK-SPIRV: OpConvertUToF
+
+;; kernel void testUToF(global uint2 *a, global float2 *res) {
+;;   res[0] = convert_float2_rtz(*a);
+;; }
+
+define dso_local spir_kernel void @testUToF(<2 x i32> addrspace(1)* nocapture noundef readonly %a, <2 x float> addrspace(1)* nocapture noundef writeonly %res) local_unnamed_addr {
+entry:
+  %0 = load <2 x i32>, <2 x i32> addrspace(1)* %a, align 8
+  %call = call spir_func <2 x float> @_Z18convert_float2_rtzDv2_j(<2 x i32> noundef %0)
+  store <2 x float> %call, <2 x float> addrspace(1)* %res, align 8
+  ret void
+}
+
+declare spir_func <2 x float> @_Z18convert_float2_rtzDv2_j(<2 x i32> noundef) local_unnamed_addr
+
+; CHECK-SPIRV: OpConvertFToU
+
+;; kernel void testFToUSat(global float2 *a, global uint2 *res) {
+;;   res[0] = convert_uint2_sat_rtn(*a);
+;; }
+
+define dso_local spir_kernel void @testFToUSat(<2 x float> addrspace(1)* nocapture noundef readonly %a, <2 x i32> addrspace(1)* nocapture noundef writeonly %res) local_unnamed_addr {
+entry:
+  %0 = load <2 x float>, <2 x float> addrspace(1)* %a, align 8
+  %call = call spir_func <2 x i32> @_Z21convert_uint2_sat_rtnDv2_f(<2 x float> noundef %0)
+  store <2 x i32> %call, <2 x i32> addrspace(1)* %res, align 8
+  ret void
+}
+
+declare spir_func <2 x i32> @_Z21convert_uint2_sat_rtnDv2_f(<2 x float> noundef) local_unnamed_addr
+
+; CHECK-SPIRV: OpSatConvertSToU
+
+;; kernel void testUToUSat(global uchar *a, global uint *res) {
+;;   res[0] = convert_uint_sat(*a);
+;; }
+
+define dso_local spir_kernel void @testUToUSat(i8 addrspace(1)* nocapture noundef readonly %a, i32 addrspace(1)* nocapture noundef writeonly %res) local_unnamed_addr {
+entry:
+  %0 = load i8, i8 addrspace(1)* %a, align 1
+  %call = call spir_func i32 @_Z16convert_uint_sath(i8 noundef zeroext %0)
+  store i32 %call, i32 addrspace(1)* %res, align 4
+  ret void
+}
+
+declare spir_func i32 @_Z16convert_uint_sath(i8 noundef zeroext) local_unnamed_addr
+
+; CHECK-SPIRV: OpSatConvertSToU
+
+;; kernel void testUToUSat1(global uint *a, global uchar *res) {
+;;   res[0] = convert_uchar_sat(*a);
+;; }
+
+define dso_local spir_kernel void @testUToUSat1(i32 addrspace(1)* nocapture noundef readonly %a, i8 addrspace(1)* nocapture noundef writeonly %res) local_unnamed_addr {
+entry:
+  %0 = load i32, i32 addrspace(1)* %a, align 4
+  %call = call spir_func zeroext i8 @_Z17convert_uchar_satj(i32 noundef %0)
+  store i8 %call, i8 addrspace(1)* %res, align 1
+  ret void
+}
+
+declare spir_func zeroext i8 @_Z17convert_uchar_satj(i32 noundef) local_unnamed_addr
+
+; CHECK-SPIRV: OpConvertFToU
+
+;; kernel void testFToU(global float3 *a, global uint3 *res) {
+;;   res[0] = convert_uint3_rtp(*a);
+;; }
+
+define dso_local spir_kernel void @testFToU(<3 x float> addrspace(1)* nocapture noundef readonly %a, <3 x i32> addrspace(1)* nocapture noundef writeonly %res) local_unnamed_addr {
+entry:
+  %castToVec4 = bitcast <3 x float> addrspace(1)* %a to <4 x float> addrspace(1)*
+  %loadVec4 = load <4 x float>, <4 x float> addrspace(1)* %castToVec4, align 16
+  %extractVec = shufflevector <4 x float> %loadVec4, <4 x float> poison, <3 x i32> <i32 0, i32 1, i32 2>
+  %call = call spir_func <3 x i32> @_Z17convert_uint3_rtpDv3_f(<3 x float> noundef %extractVec)
+  %extractVec1 = shufflevector <3 x i32> %call, <3 x i32> poison, <4 x i32> <i32 0, i32 1, i32 2, i32 undef>
+  %storetmp = bitcast <3 x i32> addrspace(1)* %res to <4 x i32> addrspace(1)*
+  store <4 x i32> %extractVec1, <4 x i32> addrspace(1)* %storetmp, align 16
+  ret void
+}
+
+declare spir_func <3 x i32> @_Z17convert_uint3_rtpDv3_f(<3 x float> noundef) local_unnamed_addr

diff  --git a/llvm/test/CodeGen/SPIRV/transcoding/image_get_size_with_access_qualifiers.ll b/llvm/test/CodeGen/SPIRV/transcoding/image_get_size_with_access_qualifiers.ll
new file mode 100644
index 0000000000000..6ef311c1d2f57
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/image_get_size_with_access_qualifiers.ll
@@ -0,0 +1,18 @@
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV-DAG: %[[#IntTyID:]] = OpTypeInt
+; CHECK-SPIRV-DAG: %[[#VoidTyID:]] = OpTypeVoid
+; CHECK-SPIRV-DAG: %[[#ImageTyID:]] = OpTypeImage %[[#VoidTyID]] 2D 0 1 0 0 Unknown ReadOnly
+; CHECK-SPIRV-DAG: %[[#VectorTyID:]] = OpTypeVector %[[#IntTyID]] [[#]]
+; CHECK-SPIRV:     %[[#ImageArgID:]] = OpFunctionParameter %[[#ImageTyID]]
+; CHECK-SPIRV:     %[[#]] = OpImageQuerySizeLod %[[#VectorTyID]] %[[#ImageArgID]]
+
+%opencl.image2d_array_ro_t = type opaque
+
+define spir_kernel void @sample_kernel(%opencl.image2d_array_ro_t addrspace(1)* %input) {
+entry:
+  %call = call spir_func i32 @_Z15get_image_width20ocl_image2d_array_ro(%opencl.image2d_array_ro_t addrspace(1)* %input)
+  ret void
+}
+
+declare spir_func i32 @_Z15get_image_width20ocl_image2d_array_ro(%opencl.image2d_array_ro_t addrspace(1)*)

diff  --git a/llvm/test/CodeGen/SPIRV/transcoding/image_with_access_qualifiers.ll b/llvm/test/CodeGen/SPIRV/transcoding/image_with_access_qualifiers.ll
new file mode 100644
index 0000000000000..00c7d61f91998
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/image_with_access_qualifiers.ll
@@ -0,0 +1,23 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV-DAG: OpCapability ImageReadWrite
+; CHECK-SPIRV-DAG: OpCapability LiteralSampler
+; CHECK-SPIRV-DAG: %[[#TyVoid:]] = OpTypeVoid
+; CHECK-SPIRV-DAG: %[[#TyImageID:]] = OpTypeImage %[[#TyVoid]] 1D 0 0 0 0 Unknown ReadWrite
+; CHECK-SPIRV-DAG: %[[#TySampledImageID:]] = OpTypeSampledImage %[[#TyImageID]]
+
+; CHECK-SPIRV-DAG: %[[#ResID:]] = OpSampledImage %[[#TySampledImageID]]
+; CHECK-SPIRV:     %[[#]] = OpImageSampleExplicitLod %[[#]] %[[#ResID]]
+
+%opencl.image1d_rw_t = type opaque
+
+define spir_func void @sampFun(%opencl.image1d_rw_t addrspace(1)* %image) {
+entry:
+  %image.addr = alloca %opencl.image1d_rw_t addrspace(1)*, align 4
+  store %opencl.image1d_rw_t addrspace(1)* %image, %opencl.image1d_rw_t addrspace(1)** %image.addr, align 4
+  %0 = load %opencl.image1d_rw_t addrspace(1)*, %opencl.image1d_rw_t addrspace(1)** %image.addr, align 4
+  %call = call spir_func <4 x float> @_Z11read_imagef14ocl_image1d_rw11ocl_sampleri(%opencl.image1d_rw_t addrspace(1)* %0, i32 8, i32 2)
+  ret void
+}
+
+declare spir_func <4 x float> @_Z11read_imagef14ocl_image1d_rw11ocl_sampleri(%opencl.image1d_rw_t addrspace(1)*, i32, i32)

diff  --git a/llvm/test/CodeGen/SPIRV/transcoding/optional-core-features-multiple.ll b/llvm/test/CodeGen/SPIRV/transcoding/optional-core-features-multiple.ll
new file mode 100644
index 0000000000000..0c666820f7938
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/optional-core-features-multiple.ll
@@ -0,0 +1,19 @@
+;; OpenCL C source
+;; -----------------------------------------------
+;; double d = 1.0;
+;; kernel void test(read_only image2d_t img) {}
+;; -----------------------------------------------
+
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+%opencl.image2d_t = type opaque
+
+ at d = addrspace(1) global double 1.000000e+00, align 8
+
+define spir_kernel void @test(%opencl.image2d_t addrspace(1)* nocapture %img) {
+entry:
+  ret void
+}
+
+; CHECK-SPIRV-DAG: OpCapability Float64
+; CHECK-SPIRV-DAG: OpCapability ImageBasic

diff  --git a/llvm/test/CodeGen/SPIRV/transcoding/spirv-types.ll b/llvm/test/CodeGen/SPIRV/transcoding/spirv-types.ll
new file mode 100644
index 0000000000000..4372273c8a4e7
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/spirv-types.ll
@@ -0,0 +1,102 @@
+;; Test SPIR-V opaque types
+
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV-DAG: OpCapability Float16
+; CHECK-SPIRV-DAG: OpCapability ImageReadWrite
+; CHECK-SPIRV-DAG: OpCapability Pipes
+; CHECK-SPIRV-DAG: OpCapability DeviceEnqueue
+
+; CHECK-SPIRV-DAG: %[[#VOID:]] = OpTypeVoid
+; CHECK-SPIRV-DAG: %[[#INT:]] = OpTypeInt 32 0
+; CHECK-SPIRV-DAG: %[[#HALF:]] = OpTypeFloat 16
+; CHECK-SPIRV-DAG: %[[#FLOAT:]] = OpTypeFloat 32
+; CHECK-SPIRV-DAG: %[[#PIPE_RD:]] = OpTypePipe ReadOnly
+; CHECK-SPIRV-DAG: %[[#PIPE_WR:]] = OpTypePipe WriteOnly
+; CHECK-SPIRV-DAG: %[[#IMG1D_RD:]] = OpTypeImage %[[#VOID]] 1D 0 0 0 0 Unknown ReadOnly
+; CHECK-SPIRV-DAG: %[[#IMG2D_RD:]] = OpTypeImage %[[#INT]] 2D 0 0 0 0 Unknown ReadOnly
+; CHECK-SPIRV-DAG: %[[#IMG3D_RD:]] = OpTypeImage %[[#INT]] 3D 0 0 0 0 Unknown ReadOnly
+; CHECK-SPIRV-DAG: %[[#IMG2DD_RD:]] = OpTypeImage %[[#FLOAT]] 2D 1 0 0 0 Unknown ReadOnly
+; CHECK-SPIRV-DAG: %[[#IMG2DA_RD:]] = OpTypeImage %[[#HALF]] 2D 0 1 0 0 Unknown ReadOnly
+; CHECK-SPIRV-DAG: %[[#IMG1DB_RD:]] = OpTypeImage %[[#FLOAT]] Buffer 0 0 0 0 Unknown ReadOnly
+; CHECK-SPIRV-DAG: %[[#IMG1D_WR:]] = OpTypeImage %[[#VOID]] 1D 0 0 0 0 Unknown WriteOnly
+; CHECK-SPIRV-DAG: %[[#IMG2D_RW:]] = OpTypeImage %[[#VOID]] 2D 0 0 0 0 Unknown ReadWrite
+; CHECK-SPIRV-DAG: %[[#DEVEVENT:]] = OpTypeDeviceEvent
+; CHECK-SPIRV-DAG: %[[#EVENT:]] = OpTypeEvent
+; CHECK-SPIRV-DAG: %[[#QUEUE:]] = OpTypeQueue
+; CHECK-SPIRV-DAG: %[[#RESID:]] = OpTypeReserveId
+; CHECK-SPIRV-DAG: %[[#SAMP:]] = OpTypeSampler
+; CHECK-SPIRV-DAG: %[[#SAMPIMG:]] = OpTypeSampledImage %[[#IMG2DD_RD]]
+
+%spirv.Pipe._0 = type opaque ; read_only pipe
+%spirv.Pipe._1 = type opaque ; write_only pipe
+%spirv.Image._void_0_0_0_0_0_0_0 = type opaque ; read_only image1d_ro_t
+%spirv.Image._int_1_0_0_0_0_0_0 = type opaque ; read_only image2d_ro_t
+%spirv.Image._uint_2_0_0_0_0_0_0 = type opaque ; read_only image3d_ro_t
+%spirv.Image._float_1_1_0_0_0_0_0 = type opaque; read_only image2d_depth_ro_t
+%spirv.Image._half_1_0_1_0_0_0_0 = type opaque ; read_only image2d_array_ro_t
+%spirv.Image._float_5_0_0_0_0_0_0 = type opaque ; read_only image1d_buffer_ro_t
+%spirv.Image._void_0_0_0_0_0_0_1 = type opaque ; write_only image1d_wo_t
+%spirv.Image._void_1_0_0_0_0_0_2 = type opaque ; read_write image2d_rw_t
+%spirv.DeviceEvent          = type opaque ; clk_event_t
+%spirv.Event                = type opaque ; event_t
+%spirv.Queue                = type opaque ; queue_t
+%spirv.ReserveId            = type opaque ; reserve_id_t
+%spirv.Sampler              = type opaque ; sampler_t
+%spirv.SampledImage._float_1_1_0_0_0_0_0 = type opaque
+
+; CHECK-SPIRV: OpFunction
+; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#PIPE_RD]]
+; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#PIPE_WR]]
+; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#IMG1D_RD]]
+; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#IMG2D_RD]]
+; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#IMG3D_RD]]
+; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#IMG2DA_RD]]
+; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#IMG1DB_RD]]
+; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#IMG1D_WR]]
+; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#IMG2D_RW]]
+
+define spir_kernel void @foo(
+  %spirv.Pipe._0 addrspace(1)* nocapture %a,
+  %spirv.Pipe._1 addrspace(1)* nocapture %b,
+  %spirv.Image._void_0_0_0_0_0_0_0 addrspace(1)* nocapture %c1,
+  %spirv.Image._int_1_0_0_0_0_0_0 addrspace(1)* nocapture %d1,
+  %spirv.Image._uint_2_0_0_0_0_0_0 addrspace(1)* nocapture %e1,
+  %spirv.Image._half_1_0_1_0_0_0_0 addrspace(1)* nocapture %f1,
+  %spirv.Image._float_5_0_0_0_0_0_0 addrspace(1)* nocapture %g1,
+  %spirv.Image._void_0_0_0_0_0_0_1 addrspace(1)* nocapture %c2,
+  %spirv.Image._void_1_0_0_0_0_0_2 addrspace(1)* nocapture %d3) {
+entry:
+  ret void
+}
+
+; CHECK-SPIRV: OpFunction
+; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#DEVEVENT]]
+; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#EVENT]]
+; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#QUEUE]]
+; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#RESID]]
+
+define spir_func void @bar(
+  %spirv.DeviceEvent * %a,
+  %spirv.Event * %b,
+  %spirv.Queue * %c,
+  %spirv.ReserveId * %d) {
+  ret void
+}
+
+; CHECK-SPIRV: OpFunction
+; CHECK-SPIRV: %[[#IMG_ARG:]] = OpFunctionParameter %[[#IMG2DD_RD]]
+; CHECK-SPIRV: %[[#SAMP_ARG:]] = OpFunctionParameter %[[#SAMP]]
+; CHECK-SPIRV: %[[#SAMPIMG_VAR:]] = OpSampledImage %[[#SAMPIMG]] %[[#IMG_ARG]] %[[#SAMP_ARG]]
+; CHECK-SPIRV: %[[#]] = OpImageSampleExplicitLod %[[#]] %[[#SAMPIMG_VAR]]
+
+define spir_func void @test_sampler(%spirv.Image._float_1_1_0_0_0_0_0 addrspace(1)* %srcimg.coerce,
+                                    %spirv.Sampler addrspace(1)* %s.coerce) {
+  %1 = tail call spir_func %spirv.SampledImage._float_1_1_0_0_0_0_0 addrspace(1)* @_Z20__spirv_SampledImagePU3AS1K34__spirv_Image__float_1_1_0_0_0_0_0PU3AS1K15__spirv_Sampler(%spirv.Image._float_1_1_0_0_0_0_0 addrspace(1)* %srcimg.coerce, %spirv.Sampler addrspace(1)* %s.coerce)
+  %2 = tail call spir_func <4 x float> @_Z38__spirv_ImageSampleExplicitLod_Rfloat4PU3AS120__spirv_SampledImageDv4_iif(%spirv.SampledImage._float_1_1_0_0_0_0_0 addrspace(1)* %1, <4 x i32> zeroinitializer, i32 2, float 1.000000e+00)
+  ret void
+}
+
+declare spir_func %spirv.SampledImage._float_1_1_0_0_0_0_0 addrspace(1)* @_Z20__spirv_SampledImagePU3AS1K34__spirv_Image__float_1_1_0_0_0_0_0PU3AS1K15__spirv_Sampler(%spirv.Image._float_1_1_0_0_0_0_0 addrspace(1)*, %spirv.Sampler addrspace(1)*)
+
+declare spir_func <4 x float> @_Z38__spirv_ImageSampleExplicitLod_Rfloat4PU3AS120__spirv_SampledImageDv4_iif(%spirv.SampledImage._float_1_1_0_0_0_0_0 addrspace(1)*, <4 x i32>, i32, float)