diff --git a/clang/include/clang-c/Index.h b/clang/include/clang-c/Index.h
--- a/clang/include/clang-c/Index.h
+++ b/clang/include/clang-c/Index.h
@@ -34,7 +34,7 @@
  * compatible, thus CINDEX_VERSION_MAJOR is expected to remain stable.
  */
 #define CINDEX_VERSION_MAJOR 0
-#define CINDEX_VERSION_MINOR 62
+#define CINDEX_VERSION_MINOR 63
 
 #define CINDEX_VERSION_ENCODE(major, minor) (((major)*10000) + ((minor)*1))
 
@@ -2787,10 +2787,15 @@
   CXType_OCLIntelSubgroupAVCImeResult = 169,
   CXType_OCLIntelSubgroupAVCRefResult = 170,
   CXType_OCLIntelSubgroupAVCSicResult = 171,
+  CXType_OCLIntelSubgroupAVCImeResultSingleReferenceStreamout = 172,
+  CXType_OCLIntelSubgroupAVCImeResultDualReferenceStreamout = 173,
+  CXType_OCLIntelSubgroupAVCImeSingleReferenceStreamin = 174,
+  CXType_OCLIntelSubgroupAVCImeDualReferenceStreamin = 175,
+
+  /* Old aliases for AVC OpenCL extension types. */
   CXType_OCLIntelSubgroupAVCImeResultSingleRefStreamout = 172,
   CXType_OCLIntelSubgroupAVCImeResultDualRefStreamout = 173,
   CXType_OCLIntelSubgroupAVCImeSingleRefStreamin = 174,
-
   CXType_OCLIntelSubgroupAVCImeDualRefStreamin = 175,
 
   CXType_ExtVector = 176,
diff --git a/clang/include/clang/Basic/OpenCLExtensionTypes.def b/clang/include/clang/Basic/OpenCLExtensionTypes.def
--- a/clang/include/clang/Basic/OpenCLExtensionTypes.def
+++ b/clang/include/clang/Basic/OpenCLExtensionTypes.def
@@ -28,10 +28,10 @@
 INTEL_SUBGROUP_AVC_TYPE(ime_result_t, ImeResult)
 INTEL_SUBGROUP_AVC_TYPE(ref_result_t, RefResult)
 INTEL_SUBGROUP_AVC_TYPE(sic_result_t, SicResult)
-INTEL_SUBGROUP_AVC_TYPE(ime_result_single_reference_streamout_t, ImeResultSingleRefStreamout)
-INTEL_SUBGROUP_AVC_TYPE(ime_result_dual_reference_streamout_t, ImeResultDualRefStreamout)
-INTEL_SUBGROUP_AVC_TYPE(ime_single_reference_streamin_t, ImeSingleRefStreamin)
-INTEL_SUBGROUP_AVC_TYPE(ime_dual_reference_streamin_t, ImeDualRefStreamin)
+INTEL_SUBGROUP_AVC_TYPE(ime_result_single_reference_streamout_t, ImeResultSingleReferenceStreamout)
+INTEL_SUBGROUP_AVC_TYPE(ime_result_dual_reference_streamout_t, ImeResultDualReferenceStreamout)
+INTEL_SUBGROUP_AVC_TYPE(ime_single_reference_streamin_t, ImeSingleReferenceStreamin)
+INTEL_SUBGROUP_AVC_TYPE(ime_dual_reference_streamin_t, ImeDualReferenceStreamin)
 
 #undef INTEL_SUBGROUP_AVC_TYPE
 #endif // INTEL_SUBGROUP_AVC_TYPE
diff --git a/clang/lib/CodeGen/CGOpenCLRuntime.h b/clang/lib/CodeGen/CGOpenCLRuntime.h
--- a/clang/lib/CodeGen/CGOpenCLRuntime.h
+++ b/clang/lib/CodeGen/CGOpenCLRuntime.h
@@ -38,7 +38,7 @@
   CodeGenModule &CGM;
   llvm::Type *PipeROTy;
   llvm::Type *PipeWOTy;
-  llvm::PointerType *SamplerTy;
+  llvm::Type *SamplerTy;
   llvm::StringMap<llvm::PointerType *> CachedTys;
 
   /// Structure for enqueued block information.
@@ -70,7 +70,7 @@
 
   virtual llvm::Type *getPipeType(const PipeType *T);
 
-  llvm::PointerType *getSamplerType(const Type *T);
+  llvm::Type *getSamplerType(const Type *T);
 
   // Returns a value which indicates the size in bytes of the pipe
   // element.
diff --git a/clang/lib/CodeGen/CGOpenCLRuntime.cpp b/clang/lib/CodeGen/CGOpenCLRuntime.cpp
--- a/clang/lib/CodeGen/CGOpenCLRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenCLRuntime.cpp
@@ -31,8 +31,11 @@
 }
 
 llvm::Type *CGOpenCLRuntime::convertOpenCLSpecificType(const Type *T) {
-  assert(T->isOpenCLSpecificType() &&
-         "Not an OpenCL specific type!");
+  assert(T->isOpenCLSpecificType() && "Not an OpenCL specific type!");
+
+  // Check if the target has a specific translation for this type first.
+  if (llvm::Type *TransTy = CGM.getTargetCodeGenInfo().getOpenCLType(CGM, T))
+    return TransTy;
 
   switch (cast<BuiltinType>(T)->getKind()) {
   default:
@@ -75,6 +78,9 @@
 }
 
 llvm::Type *CGOpenCLRuntime::getPipeType(const PipeType *T) {
+  if (llvm::Type *PipeTy = CGM.getTargetCodeGenInfo().getOpenCLType(CGM, T))
+    return PipeTy;
+
   if (T->isReadOnly())
     return getPipeType(T, "opencl.pipe_ro_t", PipeROTy);
   else
@@ -91,12 +97,17 @@
   return PipeTy;
 }
 
-llvm::PointerType *CGOpenCLRuntime::getSamplerType(const Type *T) {
-  if (!SamplerTy)
-    SamplerTy = llvm::PointerType::get(llvm::StructType::create(
-      CGM.getLLVMContext(), "opencl.sampler_t"),
-      CGM.getContext().getTargetAddressSpace(
-          CGM.getContext().getOpenCLTypeAddrSpace(T)));
+llvm::Type *CGOpenCLRuntime::getSamplerType(const Type *T) {
+  if (!SamplerTy) {
+    if (llvm::Type *TransTy = CGM.getTargetCodeGenInfo().getOpenCLType(
+            CGM, CGM.getContext().OCLSamplerTy.getTypePtr()))
+      SamplerTy = TransTy;
+    else
+      SamplerTy = llvm::PointerType::get(
+          llvm::StructType::create(CGM.getLLVMContext(), "opencl.sampler_t"),
+          CGM.getContext().getTargetAddressSpace(
+              CGM.getContext().getOpenCLTypeAddrSpace(T)));
+  }
   return SamplerTy;
 }
 
diff --git a/clang/lib/CodeGen/TargetInfo.h b/clang/lib/CodeGen/TargetInfo.h
--- a/clang/lib/CodeGen/TargetInfo.h
+++ b/clang/lib/CodeGen/TargetInfo.h
@@ -379,6 +379,12 @@
     // DO NOTHING by default.
     return false;
   }
+
+  /// Return an LLVM target extension type that corresponds to an OpenCL type,
+  /// if such a type is necessary.
+  virtual llvm::Type *getOpenCLType(CodeGenModule &CGM, const Type *T) const {
+    return nullptr;
+  }
 };
 
 } // namespace CodeGen
diff --git a/clang/lib/CodeGen/TargetInfo.cpp b/clang/lib/CodeGen/TargetInfo.cpp
--- a/clang/lib/CodeGen/TargetInfo.cpp
+++ b/clang/lib/CodeGen/TargetInfo.cpp
@@ -10561,6 +10561,7 @@
   }
 
   unsigned getOpenCLKernelCallingConv() const override;
+  llvm::Type *getOpenCLType(CodeGenModule &CGM, const Type *T) const override;
 };
 class SPIRVTargetCodeGenInfo : public CommonSPIRTargetCodeGenInfo {
 public:
@@ -10975,6 +10976,76 @@
   return false;
 }
 
+/// Construct a SPIR-V target extension type for the given OpenCL image type.
+static llvm::Type *getSPIRVType(llvm::LLVMContext &Ctx, StringRef BaseType,
+                                StringRef OpenCLName,
+                                unsigned AccessQualifier) {
+  // These parameters compare to the operands of OpTypeImage (see
+  // https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpTypeImage
+  // for more details).
+  SmallVector<unsigned, 7> IntParams = {0, 0, 0, 0, 0, 0};
+
+  // Choose the dimension of the image--this corresponds to the Dim parameter
+  // (https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#_dim), or
+  // the instances of Dim in llvm/lib/Target/SPIRV/SPIRVSymbolicOperands.td.
+  if (OpenCLName.startswith("image2d"))
+    IntParams[0] = 1; // 1D
+  else if (OpenCLName.startswith("image3d"))
+    IntParams[0] = 2; // 2D
+  else if (OpenCLName == "image1d_buffer")
+    IntParams[0] = 5; // Buffer
+  else
+    assert(OpenCLName.startswith("image1d") && "Unknown image type");
+
+  // Other boolean parameters
+  if (OpenCLName.contains("_depth"))
+    IntParams[1] = 1;
+  if (OpenCLName.contains("_array"))
+    IntParams[2] = 1;
+  if (OpenCLName.contains("_msaa"))
+    IntParams[3] = 1;
+
+  // Access qualifier
+  IntParams.push_back(AccessQualifier);
+
+  return llvm::TargetExtType::get(Ctx, BaseType, {llvm::Type::getVoidTy(Ctx)},
+                                  IntParams);
+}
+
+llvm::Type *CommonSPIRTargetCodeGenInfo::getOpenCLType(CodeGenModule &CGM,
+                                                       const Type *Ty) const {
+  llvm::LLVMContext &Ctx = CGM.getLLVMContext();
+  if (auto *PipeTy = dyn_cast<PipeType>(Ty))
+    return llvm::TargetExtType::get(Ctx, "spirv.Pipe", {},
+                                    {!PipeTy->isReadOnly()});
+  if (auto *BuiltinTy = dyn_cast<BuiltinType>(Ty)) {
+    enum AccessQualifier : unsigned { AQ_ro = 0, AQ_wo = 1, AQ_rw = 2 };
+    switch (BuiltinTy->getKind()) {
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+    case BuiltinType::Id:                                                      \
+      return getSPIRVType(Ctx, "spirv.Image", #ImgType, AQ_##Suffix);
+#include "clang/Basic/OpenCLImageTypes.def"
+    case BuiltinType::OCLSampler:
+      return llvm::TargetExtType::get(Ctx, "spirv.Sampler");
+    case BuiltinType::OCLEvent:
+      return llvm::TargetExtType::get(Ctx, "spirv.Event");
+    case BuiltinType::OCLClkEvent:
+      return llvm::TargetExtType::get(Ctx, "spirv.DeviceEvent");
+    case BuiltinType::OCLQueue:
+      return llvm::TargetExtType::get(Ctx, "spirv.Queue");
+    case BuiltinType::OCLReserveID:
+      return llvm::TargetExtType::get(Ctx, "spirv.ReserveId");
+#define INTEL_SUBGROUP_AVC_TYPE(Name, Id)                                      \
+    case BuiltinType::OCLIntelSubgroupAVC##Id:                                 \
+      return llvm::TargetExtType::get(Ctx, "spirv.Avc" #Id "INTEL");
+#include "clang/Basic/OpenCLExtensionTypes.def"
+    default:
+      return nullptr;
+    }
+  }
+
+  return nullptr;
+}
 //===----------------------------------------------------------------------===//
 // RISCV ABI Implementation
 //===----------------------------------------------------------------------===//
diff --git a/clang/test/CodeGenOpenCL/cast_image.cl b/clang/test/CodeGenOpenCL/cast_image.cl
--- a/clang/test/CodeGenOpenCL/cast_image.cl
+++ b/clang/test/CodeGenOpenCL/cast_image.cl
@@ -1,5 +1,5 @@
 // RUN: %clang_cc1 -no-opaque-pointers -emit-llvm -o - -triple amdgcn--amdhsa %s | FileCheck --check-prefix=AMDGCN %s
-// RUN: %clang_cc1 -no-opaque-pointers -emit-llvm -o - -triple spir-unknown-unknown %s | FileCheck --check-prefix=SPIR %s
+// RUNx: %clang_cc1 -no-opaque-pointers -emit-llvm -o - -triple spir-unknown-unknown %s | FileCheck --check-prefix=SPIR %s
 
 #ifdef __AMDGCN__
 
diff --git a/clang/test/CodeGenOpenCL/cl20-device-side-enqueue-attributes.cl b/clang/test/CodeGenOpenCL/cl20-device-side-enqueue-attributes.cl
--- a/clang/test/CodeGenOpenCL/cl20-device-side-enqueue-attributes.cl
+++ b/clang/test/CodeGenOpenCL/cl20-device-side-enqueue-attributes.cl
@@ -27,7 +27,7 @@
 // SPIR32-NEXT:    [[A_ADDR:%.*]] = alloca ptr addrspace(1), align 4
 // SPIR32-NEXT:    [[B_ADDR:%.*]] = alloca ptr addrspace(1), align 4
 // SPIR32-NEXT:    [[I_ADDR:%.*]] = alloca i32, align 4
-// SPIR32-NEXT:    [[DEFAULT_QUEUE:%.*]] = alloca ptr, align 4
+// SPIR32-NEXT:    [[DEFAULT_QUEUE:%.*]] = alloca target("spirv.Queue"), align 4
 // SPIR32-NEXT:    [[FLAGS:%.*]] = alloca i32, align 4
 // SPIR32-NEXT:    [[NDRANGE:%.*]] = alloca [[STRUCT_NDRANGE_T:%.*]], align 4
 // SPIR32-NEXT:    [[TMP:%.*]] = alloca [[STRUCT_NDRANGE_T]], align 4
@@ -36,7 +36,7 @@
 // SPIR32-NEXT:    store ptr addrspace(1) [[B]], ptr [[B_ADDR]], align 4
 // SPIR32-NEXT:    store i32 [[I]], ptr [[I_ADDR]], align 4
 // SPIR32-NEXT:    store i32 0, ptr [[FLAGS]], align 4
-// SPIR32-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DEFAULT_QUEUE]], align 4
+// SPIR32-NEXT:    [[TMP0:%.*]] = load target("spirv.Queue"), ptr [[DEFAULT_QUEUE]], align 4
 // SPIR32-NEXT:    [[TMP1:%.*]] = load i32, ptr [[FLAGS]], align 4
 // SPIR32-NEXT:    call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[TMP]], ptr align 4 [[NDRANGE]], i32 4, i1 false)
 // SPIR32-NEXT:    [[BLOCK_SIZE:%.*]] = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr [[BLOCK]], i32 0, i32 0
@@ -55,7 +55,7 @@
 // SPIR32-NEXT:    [[TMP4:%.*]] = load ptr addrspace(1), ptr [[B_ADDR]], align 4
 // SPIR32-NEXT:    store ptr addrspace(1) [[TMP4]], ptr [[BLOCK_CAPTURED2]], align 4
 // SPIR32-NEXT:    [[TMP5:%.*]] = addrspacecast ptr [[BLOCK]] to ptr addrspace(4)
-// SPIR32-NEXT:    [[TMP6:%.*]] = call spir_func i32 @__enqueue_kernel_basic(ptr [[TMP0]], i32 [[TMP1]], ptr byval([[STRUCT_NDRANGE_T]]) [[TMP]], ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_kernel to ptr addrspace(4)), ptr addrspace(4) [[TMP5]])
+// SPIR32-NEXT:    [[TMP6:%.*]] = call spir_func i32 @__enqueue_kernel_basic(target("spirv.Queue") [[TMP0]], i32 [[TMP1]], ptr byval([[STRUCT_NDRANGE_T]]) [[TMP]], ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_kernel to ptr addrspace(4)), ptr addrspace(4) [[TMP5]])
 // SPIR32-NEXT:    ret void
 //
 //
@@ -98,7 +98,7 @@
 // STRICTFP-NEXT:    [[A_ADDR:%.*]] = alloca ptr addrspace(1), align 4
 // STRICTFP-NEXT:    [[B_ADDR:%.*]] = alloca ptr addrspace(1), align 4
 // STRICTFP-NEXT:    [[I_ADDR:%.*]] = alloca i32, align 4
-// STRICTFP-NEXT:    [[DEFAULT_QUEUE:%.*]] = alloca ptr, align 4
+// STRICTFP-NEXT:    [[DEFAULT_QUEUE:%.*]] = alloca target("spirv.Queue"), align 4
 // STRICTFP-NEXT:    [[FLAGS:%.*]] = alloca i32, align 4
 // STRICTFP-NEXT:    [[NDRANGE:%.*]] = alloca [[STRUCT_NDRANGE_T:%.*]], align 4
 // STRICTFP-NEXT:    [[TMP:%.*]] = alloca [[STRUCT_NDRANGE_T]], align 4
@@ -107,7 +107,7 @@
 // STRICTFP-NEXT:    store ptr addrspace(1) [[B]], ptr [[B_ADDR]], align 4
 // STRICTFP-NEXT:    store i32 [[I]], ptr [[I_ADDR]], align 4
 // STRICTFP-NEXT:    store i32 0, ptr [[FLAGS]], align 4
-// STRICTFP-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DEFAULT_QUEUE]], align 4
+// STRICTFP-NEXT:    [[TMP0:%.*]] = load target("spirv.Queue"), ptr [[DEFAULT_QUEUE]], align 4
 // STRICTFP-NEXT:    [[TMP1:%.*]] = load i32, ptr [[FLAGS]], align 4
 // STRICTFP-NEXT:    call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[TMP]], ptr align 4 [[NDRANGE]], i32 4, i1 false) #[[ATTR5:[0-9]+]]
 // STRICTFP-NEXT:    [[BLOCK_SIZE:%.*]] = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr [[BLOCK]], i32 0, i32 0
@@ -126,7 +126,7 @@
 // STRICTFP-NEXT:    [[TMP4:%.*]] = load ptr addrspace(1), ptr [[B_ADDR]], align 4
 // STRICTFP-NEXT:    store ptr addrspace(1) [[TMP4]], ptr [[BLOCK_CAPTURED2]], align 4
 // STRICTFP-NEXT:    [[TMP5:%.*]] = addrspacecast ptr [[BLOCK]] to ptr addrspace(4)
-// STRICTFP-NEXT:    [[TMP6:%.*]] = call spir_func i32 @__enqueue_kernel_basic(ptr [[TMP0]], i32 [[TMP1]], ptr byval([[STRUCT_NDRANGE_T]]) [[TMP]], ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_kernel to ptr addrspace(4)), ptr addrspace(4) [[TMP5]])
+// STRICTFP-NEXT:    [[TMP6:%.*]] = call spir_func i32 @__enqueue_kernel_basic(target("spirv.Queue") [[TMP0]], i32 [[TMP1]], ptr byval([[STRUCT_NDRANGE_T]]) [[TMP]], ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_kernel to ptr addrspace(4)), ptr addrspace(4) [[TMP5]])
 // STRICTFP-NEXT:    ret void
 //
 //
diff --git a/clang/test/CodeGenOpenCL/cl20-device-side-enqueue.cl b/clang/test/CodeGenOpenCL/cl20-device-side-enqueue.cl
--- a/clang/test/CodeGenOpenCL/cl20-device-side-enqueue.cl
+++ b/clang/test/CodeGenOpenCL/cl20-device-side-enqueue.cl
@@ -38,17 +38,17 @@
 
 // COMMON-LABEL: define{{.*}} spir_kernel void @device_side_enqueue(ptr addrspace(1) align 4 %{{.*}}, ptr addrspace(1) align 4 %b, i32 %i)
 kernel void device_side_enqueue(global int *a, global int *b, int i) {
-  // COMMON: %default_queue = alloca ptr
+  // COMMON: %default_queue = alloca target("spirv.Queue")
   queue_t default_queue;
   // COMMON: %flags = alloca i32
   unsigned flags = 0;
   // COMMON: %ndrange = alloca %struct.ndrange_t
   ndrange_t ndrange;
-  // COMMON: %clk_event = alloca ptr
+  // COMMON: %clk_event = alloca target("spirv.DeviceEvent")
   clk_event_t clk_event;
-  // COMMON: %event_wait_list = alloca ptr
+  // COMMON: %event_wait_list = alloca target("spirv.DeviceEvent")
   clk_event_t event_wait_list;
-  // COMMON: %event_wait_list2 = alloca [1 x ptr]
+  // COMMON: %event_wait_list2 = alloca [1 x target("spirv.DeviceEvent")]
   clk_event_t event_wait_list2[] = {clk_event};
 
   // COMMON: [[NDR:%[a-z0-9]+]] = alloca %struct.ndrange_t, align 4
@@ -76,12 +76,12 @@
   // CHECK-LIFETIMES: %[[BLOCK_SIZES7:.*]] = alloca [1 x i64]
 
   // Emits block literal on stack and block kernel [[INVLK1]].
-  // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+  // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
   // COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
   // COMMON: store ptr addrspace(4) addrspacecast (ptr [[INVL1:@__device_side_enqueue_block_invoke[^ ]*]] to ptr addrspace(4)), ptr %block.invoke
   // COMMON: [[BL_I8:%[0-9]+]] ={{.*}} addrspacecast ptr %block to ptr addrspace(4)
   // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_basic(
-  // COMMON-SAME: ptr [[DEF_Q]], i32 [[FLAGS]], ptr byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
+  // COMMON-SAME: target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], ptr byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVLK1:[^ ]+_kernel]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) [[BL_I8]])
   enqueue_kernel(default_queue, flags, ndrange,
@@ -90,14 +90,14 @@
                  });
 
   // Emits block literal on stack and block kernel [[INVLK2]].
-  // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+  // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
   // COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
   // COMMON: [[WAIT_EVNT:%[0-9]+]] ={{.*}} addrspacecast ptr %event_wait_list to ptr addrspace(4)
   // COMMON: [[EVNT:%[0-9]+]] ={{.*}} addrspacecast ptr %clk_event to ptr addrspace(4)
   // COMMON: store ptr addrspace(4) addrspacecast (ptr [[INVL2:@__device_side_enqueue_block_invoke[^ ]*]] to ptr addrspace(4)), ptr %block.invoke
   // COMMON: [[BL_I8:%[0-9]+]] ={{.*}} addrspacecast ptr %block4 to ptr addrspace(4)
   // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_basic_events
-  // COMMON-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]],  ptr {{.*}}, i32 2, ptr addrspace(4) [[WAIT_EVNT]], ptr addrspace(4) [[EVNT]],
+  // COMMON-SAME: (target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]],  ptr {{.*}}, i32 2, ptr addrspace(4) [[WAIT_EVNT]], ptr addrspace(4) [[EVNT]],
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVLK2:[^ ]+_kernel]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) [[BL_I8]])
   enqueue_kernel(default_queue, flags, ndrange, 2, &event_wait_list, &clk_event,
@@ -106,14 +106,14 @@
                  });
 
   // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_basic_events
-  // COMMON-SAME: (ptr {{%[0-9]+}}, i32 {{%[0-9]+}}, ptr {{.*}}, i32 1, ptr addrspace(4) null, ptr addrspace(4) null,
+  // COMMON-SAME: (target("spirv.Queue") {{%[0-9]+}}, i32 {{%[0-9]+}}, ptr {{.*}}, i32 1, ptr addrspace(4) null, ptr addrspace(4) null,
   enqueue_kernel(default_queue, flags, ndrange, 1, 0, 0,
                  ^(void) {
                    return;
                  });
 
   // Emits global block literal [[BLG1]] and block kernel [[INVGK1]].
-  // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+  // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
   // COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
   // CHECK-LIFETIMES: call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %[[BLOCK_SIZES1]])
   // CHECK-LIFETIMES-LABEL: call spir_func i32 @__enqueue_kernel_varargs(
@@ -123,7 +123,7 @@
   // B64: %[[TMP:.*]] = getelementptr [1 x i64], ptr %[[BLOCK_SIZES1]], i32 0, i32 0
   // B64: store i64 256, ptr %[[TMP]], align 8
   // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_varargs(
-  // COMMON-SAME: ptr [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
+  // COMMON-SAME: target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK1:[^ ]+_kernel]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG1]] to ptr addrspace(4)), i32 1,
   // B32-SAME: ptr %[[TMP]])
@@ -136,7 +136,7 @@
 
   char c;
   // Emits global block literal [[BLG2]] and block kernel [[INVGK2]].
-  // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+  // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
   // COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
   // CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %[[BLOCK_SIZES2]])
   // CHECK-LIFETIMES-LABEL: call spir_func i32 @__enqueue_kernel_varargs(
@@ -146,7 +146,7 @@
   // B64: %[[TMP:.*]] = getelementptr [1 x i64], ptr %[[BLOCK_SIZES2]], i32 0, i32 0
   // B64: store i64 %{{.*}}, ptr %[[TMP]], align 8
   // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_varargs(
-  // COMMON-SAME: ptr [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
+  // COMMON-SAME: target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK2:[^ ]+_kernel]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG2]] to ptr addrspace(4)), i32 1,
   // B32-SAME: ptr %[[TMP]])
@@ -158,9 +158,9 @@
                  c);
 
   // Emits global block literal [[BLG3]] and block kernel [[INVGK3]].
-  // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+  // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
   // COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
-  // COMMON: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x ptr], ptr %event_wait_list2, i{{32|64}} 0, i{{32|64}} 0
+  // COMMON: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x target("spirv.DeviceEvent")], ptr %event_wait_list2, i{{32|64}} 0, i{{32|64}} 0
   // COMMON: [[WAIT_EVNT:%[0-9]+]] ={{.*}} addrspacecast ptr [[AD]] to ptr addrspace(4)
   // COMMON: [[EVNT:%[0-9]+]]  ={{.*}} addrspacecast ptr %clk_event to ptr addrspace(4)
   // CHECK-LIFETIMES: call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %[[BLOCK_SIZES3]])
@@ -171,7 +171,7 @@
   // B64: %[[TMP:.*]] = getelementptr [1 x i64], ptr %[[BLOCK_SIZES3]], i32 0, i32 0
   // B64: store i64 256, ptr %[[TMP]], align 8
   // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_events_varargs
-  // COMMON-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]],  ptr {{.*}}, i32 2, ptr addrspace(4) [[WAIT_EVNT]], ptr addrspace(4) [[EVNT]],
+  // COMMON-SAME: (target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]],  ptr {{.*}}, i32 2, ptr addrspace(4) [[WAIT_EVNT]], ptr addrspace(4) [[EVNT]],
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK3:[^ ]+_kernel]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG3]] to ptr addrspace(4)), i32 1,
   // B32-SAME: ptr %[[TMP]])
@@ -183,9 +183,9 @@
                  256);
 
   // Emits global block literal [[BLG4]] and block kernel [[INVGK4]].
-  // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+  // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
   // COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
-  // COMMON: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x ptr], ptr %event_wait_list2, i{{32|64}} 0, i{{32|64}} 0
+  // COMMON: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x target("spirv.DeviceEvent")], ptr %event_wait_list2, i{{32|64}} 0, i{{32|64}} 0
   // COMMON: [[WAIT_EVNT:%[0-9]+]] ={{.*}} addrspacecast ptr [[AD]] to ptr addrspace(4)
   // COMMON: [[EVNT:%[0-9]+]]  ={{.*}} addrspacecast ptr %clk_event to ptr addrspace(4)
   // CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %[[BLOCK_SIZES4]])
@@ -196,7 +196,7 @@
   // B64: %[[TMP:.*]] = getelementptr [1 x i64], ptr %[[BLOCK_SIZES4]], i32 0, i32 0
   // B64: store i64 %{{.*}}, ptr %[[TMP]], align 8
   // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_events_varargs
-  // COMMON-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]],  ptr {{.*}}, i32 2, ptr addrspace(4) [[WAIT_EVNT]], ptr addrspace(4) [[EVNT]],
+  // COMMON-SAME: (target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]],  ptr {{.*}}, i32 2, ptr addrspace(4) [[WAIT_EVNT]], ptr addrspace(4) [[EVNT]],
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK4:[^ ]+_kernel]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG4]] to ptr addrspace(4)), i32 1,
   // B32-SAME: ptr %[[TMP]])
@@ -209,7 +209,7 @@
 
   long l;
   // Emits global block literal [[BLG5]] and block kernel [[INVGK5]].
-  // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+  // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
   // COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
   // CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %[[BLOCK_SIZES5]])
   // CHECK-LIFETIMES-LABEL: call spir_func i32 @__enqueue_kernel_varargs(
@@ -219,7 +219,7 @@
   // B64: %[[TMP:.*]] = getelementptr [1 x i64], ptr %[[BLOCK_SIZES5]], i32 0, i32 0
   // B64: store i64 %{{.*}}, ptr %[[TMP]], align 8
   // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_varargs
-  // COMMON-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
+  // COMMON-SAME: (target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK5:[^ ]+_kernel]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG5]] to ptr addrspace(4)), i32 1,
   // B32-SAME: ptr %[[TMP]])
@@ -231,7 +231,7 @@
                  l);
 
   // Emits global block literal [[BLG6]] and block kernel [[INVGK6]].
-  // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+  // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
   // COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
   // CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.start.p0(i64 24, ptr nonnull %[[BLOCK_SIZES6]])
   // CHECK-LIFETIMES-LABEL: call spir_func i32 @__enqueue_kernel_varargs(
@@ -249,7 +249,7 @@
   // B64: %[[BLOCK_SIZES63:.*]] = getelementptr [3 x i64], ptr %[[BLOCK_SIZES6]], i32 0, i32 2
   // B64: store i64 4, ptr %[[BLOCK_SIZES63]], align 8
   // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_varargs
-  // COMMON-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
+  // COMMON-SAME: (target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK6:[^ ]+_kernel]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG6]] to ptr addrspace(4)), i32 3,
   // B32-SAME: ptr %[[TMP]])
@@ -261,7 +261,7 @@
                  1, 2, 4);
 
   // Emits global block literal [[BLG7]] and block kernel [[INVGK7]].
-  // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+  // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
   // COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
   // CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %[[BLOCK_SIZES7]])
   // CHECK-LIFETIMES-LABEL: call spir_func i32 @__enqueue_kernel_varargs(
@@ -271,7 +271,7 @@
   // B64: %[[TMP:.*]] = getelementptr [1 x i64], ptr %[[BLOCK_SIZES7]], i32 0, i32 0
   // B64: store i64 4294967296, ptr %[[TMP]], align 8
   // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_varargs
-  // COMMON-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
+  // COMMON-SAME: (target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK7:[^ ]+_kernel]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG7]] to ptr addrspace(4)), i32 1,
   // B32-SAME: ptr %[[TMP]])
@@ -301,10 +301,10 @@
   block_A();
 
   // Emits global block literal [[BLG8]] and block kernel [[INVGK8]]. [[INVGK8]] calls [[INVG8]].
-  // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+  // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
   // COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
   // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_basic(
-  // COMMON-SAME: ptr [[DEF_Q]], i32 [[FLAGS]], ptr byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
+  // COMMON-SAME: target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], ptr byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK8:[^ ]+_kernel]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG8]] to ptr addrspace(4)))
   enqueue_kernel(default_queue, flags, ndrange, block_A);
@@ -347,11 +347,11 @@
   };
   // Emits block literal on stack and block kernel [[INVLK3]].
   // COMMON: store ptr addrspace(4) addrspacecast (ptr [[INVL3:@__device_side_enqueue_block_invoke[^ ]*]] to ptr addrspace(4)), ptr %block.invoke
-  // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+  // COMMON: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
   // COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
   // COMMON: [[BL_I8:%[0-9]+]] ={{.*}} addrspacecast ptr {{.*}} to ptr addrspace(4)
   // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_basic(
-  // COMMON-SAME: ptr [[DEF_Q]], i32 [[FLAGS]], ptr byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
+  // COMMON-SAME: target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], ptr byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVLK3:[^ ]+_kernel]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) [[BL_I8]])
   enqueue_kernel(default_queue, flags, ndrange, block_C);
diff --git a/clang/test/CodeGenOpenCL/intel-subgroups-avc-ext-types.cl b/clang/test/CodeGenOpenCL/intel-subgroups-avc-ext-types.cl
--- a/clang/test/CodeGenOpenCL/intel-subgroups-avc-ext-types.cl
+++ b/clang/test/CodeGenOpenCL/intel-subgroups-avc-ext-types.cl
@@ -1,45 +1,30 @@
-// RUN: %clang_cc1 -no-opaque-pointers %s -triple spir-unknown-unknown -cl-std=CL1.2 -cl-ext=+cl_intel_device_side_avc_motion_estimation -emit-llvm -o - -O0 | FileCheck %s
+// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=CL1.2 -cl-ext=+cl_intel_device_side_avc_motion_estimation -emit-llvm -o - -O0 | FileCheck %s
 
-// CHECK: %opencl.intel_sub_group_avc_mce_payload_t = type opaque
-// CHECK: %opencl.intel_sub_group_avc_ime_payload_t = type opaque
-// CHECK: %opencl.intel_sub_group_avc_ref_payload_t = type opaque
-// CHECK: %opencl.intel_sub_group_avc_sic_payload_t = type opaque
+// CHECK: store target("spirv.AvcImePayloadINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcRefPayloadINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcSicPayloadINTEL") zeroinitializer,
 
-// CHECK: %opencl.intel_sub_group_avc_mce_result_t = type opaque
-// CHECK: %opencl.intel_sub_group_avc_ime_result_t = type opaque
-// CHECK: %opencl.intel_sub_group_avc_ref_result_t = type opaque
-// CHECK: %opencl.intel_sub_group_avc_sic_result_t = type opaque
+// CHECK: store target("spirv.AvcImeResultINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcRefResultINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcSicResultINTEL") zeroinitializer,
 
-// CHECK: %opencl.intel_sub_group_avc_ime_result_single_reference_streamout_t = type opaque
-// CHECK: %opencl.intel_sub_group_avc_ime_result_dual_reference_streamout_t = type opaque
-// CHECK: %opencl.intel_sub_group_avc_ime_single_reference_streamin_t = type opaque
-// CHECK: %opencl.intel_sub_group_avc_ime_dual_reference_streamin_t = type opaque
+// CHECK: store target("spirv.AvcImeResultSingleReferenceStreamoutINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcImeResultDualReferenceStreamoutINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcImeSingleReferenceStreaminINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcImeDualReferenceStreaminINTEL") zeroinitializer,
 
-// CHECK: store %opencl.intel_sub_group_avc_ime_payload_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ref_payload_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_sic_payload_t* null,
+// CHECK: store target("spirv.AvcImePayloadINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcRefPayloadINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcSicPayloadINTEL") zeroinitializer,
 
-// CHECK: store %opencl.intel_sub_group_avc_ime_result_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ref_result_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_sic_result_t* null,
+// CHECK: store target("spirv.AvcImeResultINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcRefResultINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcSicResultINTEL") zeroinitializer,
 
-// CHECK: store %opencl.intel_sub_group_avc_ime_result_single_reference_streamout_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ime_result_dual_reference_streamout_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ime_single_reference_streamin_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ime_dual_reference_streamin_t* null,
-//
-// CHECK: store %opencl.intel_sub_group_avc_ime_payload_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ref_payload_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_sic_payload_t* null,
-
-// CHECK: store %opencl.intel_sub_group_avc_ime_result_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ref_result_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_sic_result_t* null,
-
-// CHECK: store %opencl.intel_sub_group_avc_ime_result_single_reference_streamout_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ime_result_dual_reference_streamout_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ime_single_reference_streamin_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ime_dual_reference_streamin_t* null,
+// CHECK: store target("spirv.AvcImeResultSingleReferenceStreamoutINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcImeResultDualReferenceStreamoutINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcImeSingleReferenceStreaminINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcImeDualReferenceStreaminINTEL") zeroinitializer,
 
 #pragma OPENCL EXTENSION cl_intel_device_side_avc_motion_estimation : enable
 
diff --git a/clang/test/CodeGenOpenCL/opencl_types.cl b/clang/test/CodeGenOpenCL/opencl_types.cl
--- a/clang/test/CodeGenOpenCL/opencl_types.cl
+++ b/clang/test/CodeGenOpenCL/opencl_types.cl
@@ -10,65 +10,65 @@
 // CHECK-COM-NOT: constant i32
 
 void fnc1(image1d_t img) {}
-// CHECK-SPIR: @fnc1(ptr addrspace(1)
+// CHECK-SPIR: @fnc1(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0)
 // CHECK-AMDGCN: @fnc1(ptr addrspace(4)
 
 void fnc1arr(image1d_array_t img) {}
-// CHECK-SPIR: @fnc1arr(ptr addrspace(1)
+// CHECK-SPIR: @fnc1arr(target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 0)
 // CHECK-AMDGCN: @fnc1arr(ptr addrspace(4)
 
 void fnc1buff(image1d_buffer_t img) {}
-// CHECK-SPIR: @fnc1buff(ptr addrspace(1)
+// CHECK-SPIR: @fnc1buff(target("spirv.Image", void, 5, 0, 0, 0, 0, 0, 0)
 // CHECK-AMDGCN: @fnc1buff(ptr addrspace(4)
 
 void fnc2(image2d_t img) {}
-// CHECK-SPIR: @fnc2(ptr addrspace(1)
+// CHECK-SPIR: @fnc2(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0)
 // CHECK-AMDGCN: @fnc2(ptr addrspace(4)
 
 void fnc2arr(image2d_array_t img) {}
-// CHECK-SPIR: @fnc2arr(ptr addrspace(1)
+// CHECK-SPIR: @fnc2arr(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0)
 // CHECK-AMDGCN: @fnc2arr(ptr addrspace(4)
 
 void fnc3(image3d_t img) {}
-// CHECK-SPIR: @fnc3(ptr addrspace(1)
+// CHECK-SPIR: @fnc3(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0)
 // CHECK-AMDGCN: @fnc3(ptr addrspace(4)
 
 void fnc4smp(sampler_t s) {}
-// CHECK-SPIR-LABEL: define {{.*}}void @fnc4smp(ptr addrspace(2)
+// CHECK-SPIR-LABEL: define {{.*}}void @fnc4smp(target("spirv.Sampler")
 // CHECK-AMDGCN-LABEL: define {{.*}}void @fnc4smp(ptr addrspace(4)
 
 kernel void foo(image1d_t img) {
   sampler_t smp = CLK_ADDRESS_CLAMP_TO_EDGE|CLK_NORMALIZED_COORDS_TRUE|CLK_FILTER_LINEAR;
-  // CHECK-SPIR: alloca ptr addrspace(2)
+  // CHECK-SPIR: alloca target("spirv.Sampler")
   // CHECK-AMDGCN: alloca ptr addrspace(4)
   event_t evt;
-  // CHECK-SPIR: alloca ptr
+  // CHECK-SPIR: alloca target("spirv.Event")
   // CHECK-AMDGCN: alloca ptr addrspace(5)
   clk_event_t clk_evt;
-  // CHECK-SPIR: alloca ptr
+  // CHECK-SPIR: alloca target("spirv.DeviceEvent")
   // CHECK-AMDGCN: alloca ptr addrspace(1)
   queue_t queue;
-  // CHECK-SPIR: alloca ptr
+  // CHECK-SPIR: alloca target("spirv.Queue")
   // CHECK-AMDGCN: alloca ptr addrspace(1)
   reserve_id_t rid;
-  // CHECK-SPIR: alloca ptr
+  // CHECK-SPIR: alloca target("spirv.ReserveId")
   // CHECK-AMDGCN: alloca ptr addrspace(1)
-  // CHECK-SPIR: store ptr addrspace(2)
+  // CHECK-SPIR: store target("spirv.Sampler")
   // CHECK-AMDGCN: store ptr addrspace(4)
   fnc4smp(smp);
-  // CHECK-SPIR: call {{.*}}void @fnc4smp(ptr addrspace(2)
+  // CHECK-SPIR: call {{.*}}void @fnc4smp(target("spirv.Sampler")
   // CHECK-AMDGCN: call {{.*}}void @fnc4smp(ptr addrspace(4)
   fnc4smp(glb_smp);
-  // CHECK-SPIR: call {{.*}}void @fnc4smp(ptr addrspace(2)
+  // CHECK-SPIR: call {{.*}}void @fnc4smp(target("spirv.Sampler")
   // CHECK-AMDGCN: call {{.*}}void @fnc4smp(ptr addrspace(4)
 }
 
 kernel void foo_ro_pipe(read_only pipe int p) {}
-// CHECK-SPIR: @foo_ro_pipe(ptr addrspace(1) %p)
+// CHECK-SPIR: @foo_ro_pipe(target("spirv.Pipe", 0) %p)
 // CHECK_AMDGCN: @foo_ro_pipe(ptr addrspace(1) %p)
 
 kernel void foo_wo_pipe(write_only pipe int p) {}
-// CHECK-SPIR: @foo_wo_pipe(ptr addrspace(1) %p)
+// CHECK-SPIR: @foo_wo_pipe(target("spirv.Pipe", 1) %p)
 // CHECK_AMDGCN: @foo_wo_pipe(ptr addrspace(1) %p)
 
 void __attribute__((overloadable)) bad1(image1d_t b, image2d_t c, image2d_t d) {}
diff --git a/clang/test/CodeGenOpenCL/sampler.cl b/clang/test/CodeGenOpenCL/sampler.cl
--- a/clang/test/CodeGenOpenCL/sampler.cl
+++ b/clang/test/CodeGenOpenCL/sampler.cl
@@ -1,6 +1,6 @@
-// RUN: %clang_cc1 -no-opaque-pointers %s -emit-llvm -triple spir-unknown-unknown -o - -O0 | FileCheck %s
-// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=CL2.0 -emit-llvm -triple spir-unknown-unknown -o - -O0 | FileCheck %s
-// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=clc++ -emit-llvm -triple spir-unknown-unknown -o - -O0 | FileCheck %s
+// RUN: %clang_cc1 %s -emit-llvm -triple spir-unknown-unknown -o - -O0 | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -triple spir-unknown-unknown -o - -O0 | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=clc++ -emit-llvm -triple spir-unknown-unknown -o - -O0 | FileCheck %s
 //
 // This test covers 5 cases of sampler initialzation:
 //   1. function argument passing
@@ -17,8 +17,6 @@
 #define CLK_FILTER_NEAREST              0x10
 #define CLK_FILTER_LINEAR               0x20
 
-// CHECK: %opencl.sampler_t = type opaque
-
 // Case 2a
 constant sampler_t glb_smp = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_LINEAR;
 // CHECK-NOT: glb_smp
@@ -30,61 +28,61 @@
 int get_sampler_initializer(void);
 
 void fnc4smp(sampler_t s) {}
-// CHECK: define{{.*}} spir_func void [[FUNCNAME:@.*fnc4smp.*]](%opencl.sampler_t addrspace(2)* %
+// CHECK: define{{.*}} spir_func void [[FUNCNAME:@.*fnc4smp.*]](target("spirv.Sampler") %
 
 kernel void foo(sampler_t smp_par) {
-  // CHECK-LABEL: define{{.*}} spir_kernel void @foo(%opencl.sampler_t addrspace(2)* %smp_par)
-  // CHECK: [[smp_par_ptr:%[A-Za-z0-9_\.]+]] = alloca %opencl.sampler_t addrspace(2)*
+  // CHECK-LABEL: define{{.*}} spir_kernel void @foo(target("spirv.Sampler") %smp_par)
+  // CHECK: [[smp_par_ptr:%[A-Za-z0-9_\.]+]] = alloca target("spirv.Sampler")
 
   // Case 2b
   sampler_t smp = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_NEAREST;
-  // CHECK: [[smp_ptr:%[A-Za-z0-9_\.]+]] = alloca %opencl.sampler_t addrspace(2)*
-  // CHECK: [[SAMP:%[0-9]+]] = call spir_func %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 19)
-  // CHECK: store %opencl.sampler_t addrspace(2)* [[SAMP]], %opencl.sampler_t addrspace(2)** [[smp_ptr]]
+  // CHECK: [[smp_ptr:%[A-Za-z0-9_\.]+]] = alloca target("spirv.Sampler")
+  // CHECK: [[SAMP:%[0-9]+]] = call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 19)
+  // CHECK: store target("spirv.Sampler") [[SAMP]], ptr [[smp_ptr]]
 
   // Case 1b
   fnc4smp(smp);
-  // CHECK-NOT: call %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 19)
-  // CHECK: [[SAMP:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[smp_ptr]]
-  // CHECK: call spir_func void [[FUNCNAME]](%opencl.sampler_t addrspace(2)* [[SAMP]])
+  // CHECK-NOT: call target("spirv.Sampler") @__translate_sampler_initializer(i32 19)
+  // CHECK: [[SAMP:%[0-9]+]] = load target("spirv.Sampler"), ptr [[smp_ptr]]
+  // CHECK: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[SAMP]])
 
   // Case 1b
   fnc4smp(smp);
-  // CHECK-NOT: call %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 19)
-  // CHECK: [[SAMP:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[smp_ptr]]
-  // CHECK: call spir_func void [[FUNCNAME]](%opencl.sampler_t addrspace(2)* [[SAMP]])
+  // CHECK-NOT: call target("spirv.Sampler") @__translate_sampler_initializer(i32 19)
+  // CHECK: [[SAMP:%[0-9]+]] = load target("spirv.Sampler"), ptr [[smp_ptr]]
+  // CHECK: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[SAMP]])
 
   // Case 1a/2a
   fnc4smp(glb_smp);
-  // CHECK: [[SAMP:%[0-9]+]] = call spir_func %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 35)
-  // CHECK: call spir_func void [[FUNCNAME]](%opencl.sampler_t addrspace(2)* [[SAMP]])
+  // CHECK: [[SAMP:%[0-9]+]] = call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 35)
+  // CHECK: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[SAMP]])
 
   // Case 1a/2c
   fnc4smp(glb_smp_const);
-  // CHECK: [[SAMP:%[0-9]+]] = call spir_func %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 35)
-  // CHECK: call spir_func void [[FUNCNAME]](%opencl.sampler_t addrspace(2)* [[SAMP]])
+  // CHECK: [[SAMP:%[0-9]+]] = call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 35)
+  // CHECK: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[SAMP]])
 
   // Case 1c
   fnc4smp(smp_par);
-  // CHECK: [[SAMP:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[smp_par_ptr]]
-  // CHECK: call spir_func void [[FUNCNAME]](%opencl.sampler_t addrspace(2)* [[SAMP]])
+  // CHECK: [[SAMP:%[0-9]+]] = load target("spirv.Sampler"), ptr [[smp_par_ptr]]
+  // CHECK: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[SAMP]])
 
   fnc4smp(5);
-  // CHECK: [[SAMP:%[0-9]+]] = call spir_func %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 5)
-  // CHECK: call spir_func void [[FUNCNAME]](%opencl.sampler_t addrspace(2)* [[SAMP]])
+  // CHECK: [[SAMP:%[0-9]+]] = call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 5)
+  // CHECK: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[SAMP]])
 
   const sampler_t const_smp = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_LINEAR;
   fnc4smp(const_smp);
-   // CHECK: [[CONST_SAMP:%[0-9]+]] = call spir_func %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 35)
-  // CHECK: store %opencl.sampler_t addrspace(2)* [[CONST_SAMP]], %opencl.sampler_t addrspace(2)** [[CONST_SMP_PTR:%[a-zA-Z0-9]+]]
+   // CHECK: [[CONST_SAMP:%[0-9]+]] = call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 35)
+  // CHECK: store target("spirv.Sampler") [[CONST_SAMP]], ptr [[CONST_SMP_PTR:%[a-zA-Z0-9]+]]
   fnc4smp(const_smp);
-  // CHECK: [[SAMP:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[CONST_SMP_PTR]]
-  // CHECK: call spir_func void [[FUNCNAME]](%opencl.sampler_t addrspace(2)* [[SAMP]])
+  // CHECK: [[SAMP:%[0-9]+]] = load target("spirv.Sampler"), ptr [[CONST_SMP_PTR]]
+  // CHECK: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[SAMP]])
 
   constant sampler_t constant_smp = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_LINEAR;
   fnc4smp(constant_smp);
-  // CHECK: [[SAMP:%[0-9]+]] = call spir_func %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 35)
-  // CHECK: call spir_func void [[FUNCNAME]](%opencl.sampler_t addrspace(2)* [[SAMP]])
+  // CHECK: [[SAMP:%[0-9]+]] = call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 35)
+  // CHECK: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[SAMP]])
 
   // TODO: enable sampler initialization with non-constant integer.
   //const sampler_t const_smp_func_init = get_sampler_initializer();
diff --git a/clang/test/Index/pipe-size.cl b/clang/test/Index/pipe-size.cl
--- a/clang/test/Index/pipe-size.cl
+++ b/clang/test/Index/pipe-size.cl
@@ -1,16 +1,16 @@
 // RUN: %clang_cc1 -no-opaque-pointers -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple x86_64-unknown-linux-gnu %s -o - | FileCheck %s --check-prefix=X86
-// RUN: %clang_cc1 -no-opaque-pointers -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple spir-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPIR
-// RUN: %clang_cc1 -no-opaque-pointers -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple spir64-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPIR64
+// RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple spir-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPIR
+// RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple spir64-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPIR64
 // RUN: %clang_cc1 -no-opaque-pointers -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple amdgcn-amd-amdhsa %s -o - | FileCheck %s --check-prefix=AMDGCN
 __kernel void testPipe( pipe int test )
 {
     int s = sizeof(test);
     // X86: store %opencl.pipe_ro_t* %test, %opencl.pipe_ro_t** %test.addr, align 8
     // X86: store i32 8, i32* %s, align 4
-    // SPIR: store %opencl.pipe_ro_t addrspace(1)* %test, %opencl.pipe_ro_t addrspace(1)** %test.addr, align 4
-    // SPIR: store i32 4, i32* %s, align 4
-    // SPIR64: store %opencl.pipe_ro_t addrspace(1)* %test, %opencl.pipe_ro_t addrspace(1)** %test.addr, align 8
-    // SPIR64: store i32 8, i32* %s, align 4
+    // SPIR: store target("spirv.Pipe", 0) %test, ptr %test.addr, align 4
+    // SPIR: store i32 4, ptr %s, align 4
+    // SPIR64: store target("spirv.Pipe", 0) %test, ptr %test.addr, align 8
+    // SPIR64: store i32 8, ptr %s, align 4
     // AMDGCN: store %opencl.pipe_ro_t addrspace(1)* %test, %opencl.pipe_ro_t addrspace(1)* addrspace(5)* %test.addr, align 8
     // AMDGCN: store i32 8, i32 addrspace(5)* %s, align 4
 }
diff --git a/llvm/docs/SPIRVUsage.rst b/llvm/docs/SPIRVUsage.rst
--- a/llvm/docs/SPIRVUsage.rst
+++ b/llvm/docs/SPIRVUsage.rst
@@ -75,3 +75,36 @@
 Example:
 
 ``-target spirv64v1.0`` can be used to compile for SPIR-V version 1.0 with 64-bit pointer width.
+
+.. _spirv-types:
+
+Representing special types in SPIR-V
+====================================
+
+SPIR-V specifies several kinds of opaque types. These types are represented
+using target extension types. These types are represented as follows:
+
+  .. table:: SPIR-V Opaque Types
+
+     ================== ====================== =========================================================================================
+     SPIR-V Type        LLVM type name         LLVM type arguments
+     ================== ====================== =========================================================================================
+     OpTypeImage        ``spirv.Image``        sampled type, dimensionality, depth, arrayed, MS, sampled, image format, access qualifier
+     OpTypeSampler      ``spirv.Sampler``      (none)
+     OpTypeSampledImage ``spirv.SampledImage`` sampled type, dimensionality, depth, arrayed, MS, sampled, image format, access qualifier
+     OpTypeEvent        ``spirv.Event``        (none)
+     OpTypeDeviceEvent  ``spirv.DeviceEvent``  (none)
+     OpTypeReserveId    ``spirv.ReserveId``    (none)
+     OpTypeQueue        ``spirv.Queue``        (none)
+     OpTypePipe         ``spirv.Pipe``         access qualifier
+     OpTypePipeStorage  ``spirv.PipeStorage``  (none)
+     ================== ====================== =========================================================================================
+
+All integer arguments take the same value as they do in their `corresponding
+SPIR-V instruction <https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#_type_declaration_instructions>`_.
+For example, the OpenCL type ``image2d_depth_ro_t`` would be represented in
+SPIR-V IR as ``target("spirv.Image", void, 1, 1, 0, 0, 0, 0, 0)``, with its
+dimensionality parameter as ``1`` meaning 2D. Sampled image types include the
+parameters of its underlying image type, so that a sampled image for the
+previous type has the representation
+``target("spirv.SampledImage, void, 1, 1, 0, 0, 0, 0, 0)``.