diff --git a/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp b/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp
index 8219935fc5bec..97c76bad2eb06 100644
--- a/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp
+++ b/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp
@@ -87,7 +87,7 @@ using namespace OCLUtil;
 
 namespace SPIRV {
 
-cl::opt<bool> SPIRVMemToReg("spirv-mem2reg", cl::init(true),
+cl::opt<bool> SPIRVMemToReg("spirv-mem2reg", cl::init(false),
                             cl::desc("LLVM/SPIR-V translation enable mem2reg"));
 
 cl::opt<bool> SPIRVNoDerefAttr(
@@ -258,7 +258,7 @@ SPIRVType *LLVMToSPIRV::transType(Type *T) {
     return mapType(T, BM->addFloatType(T->getPrimitiveSizeInBits()));
 
   // A pointer to image or pipe type in LLVM is translated to a SPIRV
-  // sampler or pipe type.
+  // (non-pointer) image or pipe type.
   if (T->isPointerTy()) {
     auto ET = T->getPointerElementType();
     assert(!ET->isFunctionTy() && "Function pointer type is not allowed");
diff --git a/llvm-spirv/test/AtomicCompareExchange_cl20.ll b/llvm-spirv/test/AtomicCompareExchange_cl20.ll
index b7a16a23fb0c7..f81d052e46d10 100644
--- a/llvm-spirv/test/AtomicCompareExchange_cl20.ll
+++ b/llvm-spirv/test/AtomicCompareExchange_cl20.ll
@@ -11,9 +11,6 @@ target triple = "spir-unknown-unknown"
 ; Int64Atomics capability must be declared only if atomic builtins have 64-bit integers arguments.
 ; CHECK-NOT: Capability Int64Atomics
 
-; CHECK: Name [[Pointer:[0-9]+]] "object"
-; CHECK: Name [[ComparatorPtr:[0-9]+]] "expected"
-; CHECK: Name [[Value:[0-9]+]] "desired"
 ; CHECK: 4 TypeInt [[int:[0-9]+]] 32 0
 ; CHECK: Constant [[int]] [[DeviceScope:[0-9]+]] 1
 ; CHECK: Constant [[int]] [[SequentiallyConsistent_MS:[0-9]+]] 16
@@ -22,9 +19,9 @@ target triple = "spir-unknown-unknown"
 
 ; Function Attrs: nounwind
 define spir_func void @test(i32 addrspace(4)* %object, i32 addrspace(4)* %expected, i32 %desired) #0 {
-; CHECK: FunctionParameter [[int_ptr]] [[Pointer]]
-; CHECK: FunctionParameter [[int_ptr]] [[ComparatorPtr]]
-; CHECK: FunctionParameter [[int]] [[Value]]
+; CHECK: FunctionParameter [[int_ptr]] [[object:[0-9]+]]
+; CHECK: FunctionParameter [[int_ptr]] [[expected:[0-9]+]]
+; CHECK: FunctionParameter [[int]] [[desired:[0-9]+]]
 
 entry:
   %object.addr = alloca i32 addrspace(4)*, align 4
@@ -39,11 +36,16 @@ entry:
   %0 = load i32 addrspace(4)*, i32 addrspace(4)** %object.addr, align 4
   %1 = load i32 addrspace(4)*, i32 addrspace(4)** %expected.addr, align 4
   %2 = load i32, i32* %desired.addr, align 4
-
-  %call = call spir_func zeroext i1 @_Z30atomic_compare_exchange_strongPVU3AS4U7_AtomiciPU3AS4ii(i32 addrspace(4)* %0, i32 addrspace(4)* %1, i32 %2)
-; CHECK: Load [[int]] [[Comparator:[0-9]+]] [[ComparatorPtr]]
+; CHECK: Store [[object_addr:[0-9]+]] [[object]]
+; CHECK: Store [[expected_addr:[0-9]+]] [[expected]]
+; CHECK: Store [[desired_addr:[0-9]+]] [[desired]]
+; CHECK: Load [[int_ptr]] [[Pointer:[0-9]+]] [[object_addr]]
+; CHECK: Load [[int_ptr]] [[exp:[0-9]+]] [[expected_addr]]
+; CHECK: Load [[int]] [[Value:[0-9]+]] [[desired_addr]]
+; CHECK: Load [[int]] [[Comparator:[0-9]+]] [[exp]]
 ; CHECK-NEXT: 9 AtomicCompareExchange [[int]] [[Result:[0-9]+]] [[Pointer]] [[DeviceScope]] [[SequentiallyConsistent_MS]] [[SequentiallyConsistent_MS]] [[Value]] [[Comparator]]
-; CHECK-NEXT: Store [[ComparatorPtr]] [[Result]]
+  %call = call spir_func zeroext i1 @_Z30atomic_compare_exchange_strongPVU3AS4U7_AtomiciPU3AS4ii(i32 addrspace(4)* %0, i32 addrspace(4)* %1, i32 %2)
+; CHECK-NEXT: Store [[exp]] [[Result]]
 ; CHECK-NEXT: IEqual [[bool]] [[CallRes:[0-9]+]] [[Result]] [[Comparator]]
 ; CHECK-NOT: [[Result]]
   %frombool = zext i1 %call to i8
@@ -57,10 +59,13 @@ entry:
   %5 = load i32 addrspace(4)*, i32 addrspace(4)** %expected.addr, align 4
   %6 = load i32, i32* %desired.addr, align 4
 
+; CHECK: Load [[int_ptr]] [[Pointer:[0-9]+]] [[object_addr]]
+; CHECK: Load [[int_ptr]] [[exp:[0-9]+]] [[expected_addr]]
+; CHECK: Load [[int]] [[Value:[0-9]+]] [[desired_addr]]
+; CHECK: Load [[int]] [[ComparatorWeak:[0-9]+]] [[exp]]
   %call2 = call spir_func zeroext i1 @_Z28atomic_compare_exchange_weakPVU3AS4U7_AtomiciPU3AS4ii(i32 addrspace(4)* %4, i32 addrspace(4)* %5, i32 %6)
-; CHECK: Load [[int]] [[ComparatorWeak:[0-9]+]] [[ComparatorPtr]]
 ; CHECK-NEXT: 9 AtomicCompareExchangeWeak [[int]] [[Result:[0-9]+]] [[Pointer]] [[DeviceScope]] [[SequentiallyConsistent_MS]] [[SequentiallyConsistent_MS]] [[Value]] [[ComparatorWeak]]
-; CHECK-NEXT: Store [[ComparatorPtr]] [[Result]]
+; CHECK-NEXT: Store [[exp]] [[Result]]
 ; CHECK-NEXT: IEqual [[bool]] [[CallRes:[0-9]+]] [[Result]] [[ComparatorWeak]]
 ; CHECK-NOT: [[Result]]
 
diff --git a/llvm-spirv/test/DebugInfo/BuiltinCallLocation.cl b/llvm-spirv/test/DebugInfo/BuiltinCallLocation.cl
index 49b0cdd6fcc4c..38fc45c090252 100644
--- a/llvm-spirv/test/DebugInfo/BuiltinCallLocation.cl
+++ b/llvm-spirv/test/DebugInfo/BuiltinCallLocation.cl
@@ -8,7 +8,7 @@
 // CHECK-SPIRV: Label
 // CHECK-SPIRV: ExtInst {{.*}} DebugScope
 // CHECK-SPIRV: ExtInst {{.*}} sin
-// CHECK-LLVM: call spir_func float @_Z3sinf(float %x) {{.*}} !dbg ![[loc:[0-9]+]]
+// CHECK-LLVM: call spir_func float @_Z3sinf(float %{{.*}}) {{.*}} !dbg ![[loc:[0-9]+]]
 // CHECK-LLVM: ![[loc]] = !DILocation(line: 14, column: 10, scope: !{{.*}})
 float f(float x) {
   return sin(x);
diff --git a/llvm-spirv/test/DebugInfo/DebugDeclareUnused.cl b/llvm-spirv/test/DebugInfo/DebugDeclareUnused.cl
index f51c790aa74b8..9e624a3fac6b5 100644
--- a/llvm-spirv/test/DebugInfo/DebugDeclareUnused.cl
+++ b/llvm-spirv/test/DebugInfo/DebugDeclareUnused.cl
@@ -1,7 +1,7 @@
 // Check that we can translate llvm.dbg.declare for a local variable which was
-// deleted by mem2reg pass(enabled by default in llvm-spirv)
+// deleted by mem2reg pass(disabled by default in llvm-spirv)
 
-// RUN: %clang_cc1 %s -triple spir -disable-llvm-passes -debug-info-kind=standalone -emit-llvm-bc -o - | llvm-spirv -o %t.spv
+// RUN: %clang_cc1 %s -triple spir -disable-llvm-passes -debug-info-kind=standalone -emit-llvm-bc -o - | llvm-spirv -spirv-mem2reg -o %t.spv
 // RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
 // RUN: llvm-spirv -r %t.spv -o - | llvm-dis -o - | FileCheck %s --check-prefix=CHECK-LLVM
 
diff --git a/llvm-spirv/test/OpSwitch32.ll b/llvm-spirv/test/OpSwitch32.ll
index 33a88f2b33ec8..046b1e297eb28 100644
--- a/llvm-spirv/test/OpSwitch32.ll
+++ b/llvm-spirv/test/OpSwitch32.ll
@@ -30,7 +30,7 @@ target triple = "spir64-unknown-unknown"
 
 ;CHECK-LLVM: test_32
 ;CHECK-LLVM: entry
-;CHECK-LLVM: switch i32 %conv, label %sw.epilog
+;CHECK-LLVM: switch i32 %0, label %sw.epilog
 ;CHECK-LLVM: i32 0, label %sw.bb
 ;CHECK-LLVM: i32 1, label %sw.bb1
 
diff --git a/llvm-spirv/test/OpSwitch64.ll b/llvm-spirv/test/OpSwitch64.ll
index cc032226c9008..430d576350e93 100644
--- a/llvm-spirv/test/OpSwitch64.ll
+++ b/llvm-spirv/test/OpSwitch64.ll
@@ -33,7 +33,7 @@ target triple = "spir64-unknown-unknown"
 
 ;CHECK-LLVM: test_64
 ;CHECK-LLVM: entry
-;CHECK-LLVM: switch i64 %call, label %sw.epilog [
+;CHECK-LLVM: switch i64 %0, label %sw.epilog [
 ;CHECK-LLVM: i64 0, label %sw.bb
 ;CHECK-LLVM: i64 1, label %sw.bb1
 ;CHECK-LLVM: i64 21474836481, label %sw.bb3
diff --git a/llvm-spirv/test/image-unoptimized.cl b/llvm-spirv/test/image-unoptimized.cl
new file mode 100644
index 0000000000000..ad7bcb0fe5b95
--- /dev/null
+++ b/llvm-spirv/test/image-unoptimized.cl
@@ -0,0 +1,47 @@
+// RUN: %clang_cc1 %s -triple spir -O0 -emit-llvm-bc -o %t.bc
+// RUN: llvm-spirv %t.bc -o %t.spv
+// RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s
+// RUN: spirv-val %t.spv
+
+// CHECK: TypeImage [[TypeImage:[0-9]+]]
+// CHECK: TypeSampler [[TypeSampler:[0-9]+]]
+// CHECK: TypePointer [[TypeImagePtr:[0-9]+]] {{[0-9]+}} [[TypeImage]]
+// CHECK: TypePointer [[TypeSamplerPtr:[0-9]+]] {{[0-9]+}} [[TypeSampler]]
+
+// CHECK: FunctionParameter [[TypeImage]] [[srcimg:[0-9]+]]
+// CHECK: FunctionParameter [[TypeSampler]] [[sampler:[0-9]+]]
+
+// CHECK: Variable [[TypeImagePtr]] [[srcimg_addr:[0-9]+]]
+// CHECK: Variable [[TypeSamplerPtr]] [[sampler_addr:[0-9]+]]
+
+// CHECK: Store [[srcimg_addr]] [[srcimg]]
+// CHECK: Store [[sampler_addr]] [[sampler]]
+
+// CHECK: Load {{[0-9]+}} [[srcimg_val:[0-9]+]] [[srcimg_addr]]
+// CHECK: Load {{[0-9]+}} [[sampler_val:[0-9]+]] [[sampler_addr]]
+
+// CHECK: SampledImage {{[0-9]+}} {{[0-9]+}} [[srcimg_val]] [[sampler_val]]
+// CHECK-NEXT: ImageSampleExplicitLod
+
+// CHECK: Load {{[0-9]+}} [[srcimg_val:[0-9]+]] [[srcimg_addr]]
+// CHECK: ImageQuerySizeLod {{[0-9]+}} {{[0-9]+}} [[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});
+}
diff --git a/llvm-spirv/test/image-unoptimized.ll b/llvm-spirv/test/image-unoptimized.ll
deleted file mode 100644
index 77ebf5e199cc7..0000000000000
--- a/llvm-spirv/test/image-unoptimized.ll
+++ /dev/null
@@ -1,77 +0,0 @@
-; RUN: llvm-as %s -o %t.bc
-; RUN: llvm-spirv %t.bc -spirv-text -o %t
-; RUN: FileCheck < %t %s
-; RUN: llvm-spirv %t.bc -o %t.spv
-; RUN: spirv-val %t.spv
-target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
-target triple = "spir-unknown-unknown"
-
-%opencl.image2d_t = type opaque
-
-; Function Attrs: nounwind
-; CHECK: {{[0-9]*}} Store
-; CHECK-NEXT: 1 Return
-define spir_kernel void @test_fn(%opencl.image2d_t addrspace(1)* %srcimg, i32 %sampler, <4 x float> addrspace(1)* %results) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !4 !kernel_arg_type_qual !5 {
-entry:
-  %srcimg.addr = alloca %opencl.image2d_t addrspace(1)*, align 4
-  %sampler.addr = alloca i32, 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_t addrspace(1)* %srcimg, %opencl.image2d_t addrspace(1)** %srcimg.addr, align 4
-  store i32 %sampler, i32* %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 0) #2
-  store i32 %call, i32* %tid_x, align 4
-  %call1 = call spir_func i32 @_Z13get_global_idj(i32 1) #2
-  store i32 %call1, i32* %tid_y, align 4
-  %0 = load %opencl.image2d_t addrspace(1)*, %opencl.image2d_t addrspace(1)** %srcimg.addr, align 4
-  %1 = load i32, i32* %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
-  %4 = load <2 x i32>, <2 x i32>* %.compoundliteral
-  %call3 = call spir_func <4 x float> @_Z11read_imagef11ocl_image2d11ocl_samplerDv2_i(%opencl.image2d_t addrspace(1)* %0, i32 %1, <2 x i32> %4) #2
-  %5 = load i32, i32* %tid_y, align 4
-  %6 = load %opencl.image2d_t addrspace(1)*, %opencl.image2d_t addrspace(1)** %srcimg.addr, align 4
-  %call4 = call spir_func i32 @_Z15get_image_width11ocl_image2d(%opencl.image2d_t addrspace(1)* %6) #2
-  %mul = mul nsw i32 %5, %call4
-  %7 = load i32, i32* %tid_x, align 4
-  %add = add nsw i32 %mul, %7
-  %8 = load <4 x float> addrspace(1)*, <4 x float> addrspace(1)** %results.addr, align 4
-  %arrayidx = getelementptr inbounds <4 x float>, <4 x float> addrspace(1)* %8, i32 %add
-  store <4 x float> %call3, <4 x float> addrspace(1)* %arrayidx, align 16
-  ret void
-}
-
-; Function Attrs: nounwind readnone
-declare spir_func i32 @_Z13get_global_idj(i32) #1
-
-; Function Attrs: nounwind readnone
-declare spir_func <4 x float> @_Z11read_imagef11ocl_image2d11ocl_samplerDv2_i(%opencl.image2d_t addrspace(1)*, i32, <2 x i32>) #1
-
-; Function Attrs: nounwind readnone
-declare spir_func i32 @_Z15get_image_width11ocl_image2d(%opencl.image2d_t addrspace(1)*) #1
-
-attributes #0 = { nounwind "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-realign-stack" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
-attributes #1 = { nounwind readnone "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-realign-stack" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
-attributes #2 = { nounwind readnone }
-
-!opencl.enable.FP_CONTRACT = !{}
-!opencl.spir.version = !{!6}
-!opencl.ocl.version = !{!6}
-!opencl.used.extensions = !{!7}
-!opencl.used.optional.core.features = !{!8}
-!opencl.compiler.options = !{!7}
-
-!1 = !{i32 1, i32 0, i32 1}
-!2 = !{!"read_only", !"none", !"none"}
-!3 = !{!"image2d_t", !"sampler_t", !"float4*"}
-!4 = !{!"image2d_t", !"sampler_t", !"float4*"}
-!5 = !{!"", !"", !""}
-!6 = !{i32 1, i32 2}
-!7 = !{}
-!8 = !{!"cl_images"}