[SPIR-V] Remove -opaque-pointers=0 from LITs, fixes for opaque pointers support

[mpaszkowski]

Summary

This patch removes the -opaque-pointers=0 flag from all* the remaining tests. Some significant changes to the SPIRVGlobalRegistry, SPIRVDupliactesTracker, Intrinsics were necessary -- most notably the introduction of spv_assign_ptr_type intrinsic for tracking the underlying type of pointers.

*The following relevant tests are removed by this patch (they still require additional adjustments and fixes in the backend, will reland in subsequent patches):

• CodeGen/SPIRV/EnqueueEmptyKernel.ll
• CodeGen/SPIRV/OpenCL/basic/vstore_private.ll
• CodeGen/SPIRV/OpenCL/execute_block.ll
• CodeGen/SPIRV/half_no_extension.ll
• CodeGen/SPIRV/transcoding/BuildNDRange_2.ll
• CodeGen/SPIRV/transcoding/block_w_struct_return.ll
• CodeGen/SPIRV/transcoding/enqueue_kernel.ll
• CodeGen/SPIRV/transcoding/global_block.ll

In the upcoming week I will do best to recover the CTS testing pipeline. This requires generating of opaque IR (opaque pointers, target extension type, ...) from Intel Graphics Compiler based on LLVM 14. I expect the CTS pass rate to drop significantly.

In the meantime, I am preparing a patch dropping support for typed pointers (removing references to APIs), OpenCL/SPIR-V builtin opaque types, and any other "dead" parts of code. This patch will also partially tackle the issues discovered in the LITs from the list above and failing CTS tests.

There are still known issues which might be now easier to work on since we no longer need to support typed pointers:

• issues with lowering enqueue, ndrange builtins
• issues with lowering some structs in opaque pointer mode
• duplicate getOrCreateSPIRVPointerType methods in SPIRVGlobalRegistry
• opaque pointer handling in SPIRVDuplicatesTracker may require major redesign

I hope to discuss some of these issues during the SPIR-V Backend WG sync on Thursday.

[iliya-diyachkov]

In general it looks good. Please keep failed tests to view the progress of development and allow others to troubleshoot.

llvm/lib/Target/SPIRV/SPIRVDuplicatesTracker.h

27

Is it used inside SPIRVDuplicatesTracker?

llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp

933

Does it affect DT?

llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.h

283 ↗

(On Diff #543264)

It might be better to keep both interfaces removing duplication inside them.

llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp

1487

Does it always have a pointer type?

[bogner]

llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp

751

isOpaquePointerTy is deprecated and just returns isPointerTy as of 2ea5aa1c96cf, so the else branch here (and in other similar cases) is unreachable and can be removed.

[iliya-diyachkov]

Sorry for delay with the answer. It looks good in general.

Some tests look deleted and some look newly created, but they are just moved/renamed with many changes inside. Can we split these operations? I.e. make changes related to opaque pointers and builtin types in this patch and rename the tests in a separate patch. It will save the history of changes in the tests.

[mpaszkowski]

In D156049#4636408, @iliya-diyachkov wrote:

Sorry for delay with the answer. It looks good in general.

Some tests look deleted and some look newly created, but they are just moved/renamed with many changes inside. Can we split these operations? I.e. make changes related to opaque pointers and builtin types in this patch and rename the tests in a separate patch. It will save the history of changes in the tests.

Sure, I have made this patch a bit smaller. Still all the fixes are present, but I have kept all the LIT names the same and kept them in the same directories. Renaming, moving, and merging of tests will come in a separate patch.

[iliya-diyachkov]

Sure, I have made this patch a bit smaller. Still all the fixes are present, but I have kept all the LIT names the same and kept them in the same directories. Renaming, moving, and merging of tests will come in a separate patch.

Thanks! What about removed tests, I'm not sure if we have an analogue of linked-list.ll (with check of OpTypeOpaque). If it's true, is it possible to write such a test?

[mpaszkowski]

In D156049#4641370, @iliya-diyachkov wrote:

Sure, I have made this patch a bit smaller. Still all the fixes are present, but I have kept all the LIT names the same and kept them in the same directories. Renaming, moving, and merging of tests will come in a separate patch.

Thanks! What about removed tests, I'm not sure if we have an analogue of linked-list.ll (with check of OpTypeOpaque). If it's true, is it possible to write such a test?

Hi @iliya-diyachkov! You are right about this test, I have restored it and rebased the patch. I will work on the linked lists later. When it comes to structs, I have a patch with additional tests for which I will create an independent review. Do you think I could continue with this patch in the current form? This will allow us to address next issues in smaller patches and should unblock Nathan's changes.

[iliya-diyachkov]

Thanks Michal, it look good to me.

llvm/include/llvm/IR/IntrinsicsSPIRV.td

 //===- IntrinsicsSPIRV.td - Defines SPIRV intrinsics -------*- tablegen -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file defines all of the SPIRV-specific intrinsics.
 //
 //===----------------------------------------------------------------------===//
 
 let TargetPrefix = "spv" in {
   def int_spv_assign_type : Intrinsic<[], [llvm_any_ty, llvm_metadata_ty]>;
+  def int_spv_assign_ptr_type : Intrinsic<[], [llvm_any_ty, llvm_metadata_ty]>;
   def int_spv_assign_name : Intrinsic<[], [llvm_any_ty, llvm_vararg_ty]>;
 
   def int_spv_track_constant : Intrinsic<[llvm_any_ty], [llvm_any_ty, llvm_metadata_ty]>;
   def int_spv_init_global : Intrinsic<[], [llvm_any_ty, llvm_any_ty]>;
   def int_spv_unref_global : Intrinsic<[], [llvm_any_ty]>;
 
   def int_spv_gep : Intrinsic<[llvm_anyptr_ty], [llvm_i1_ty, llvm_any_ty, llvm_vararg_ty], [ImmArg<ArgIndex<0>>]>;
   def int_spv_load : Intrinsic<[llvm_i32_ty], [llvm_anyptr_ty, llvm_i16_ty, llvm_i8_ty], [ImmArg<ArgIndex<1>>, ImmArg<ArgIndex<2>>]>;

llvm/lib/Target/SPIRV/SPIRVDuplicatesTracker.h

 #include "llvm/ADT/DenseMap.h"
 #include "llvm/ADT/MapVector.h"
 #include "llvm/CodeGen/GlobalISel/MachineIRBuilder.h"
 #include "llvm/CodeGen/MachineModuleInfo.h"
 
 #include <type_traits>
 
 namespace llvm {
 namespace SPIRV {
 // NOTE: using MapVector instead of DenseMap because it helps getting
 // everything ordered in a stable manner for a price of extra (NumKeys)*PtrSize
 // memory and expensive removals which do not happen anyway.
 class DTSortableEntry : public MapVector<const MachineFunction *, Register> {
   SmallVector<DTSortableEntry *, 2> Deps;
 
   struct FlagsTy {
     unsigned IsFunc : 1;

 struct SpecialTypeDescriptor {
   enum SpecialTypeKind {
     STK_Empty = 0,
     STK_Image,
     STK_SampledImage,
     STK_Sampler,
     STK_Pipe,
     STK_DeviceEvent,
+    STK_Pointer,
     STK_Last = -1
   };
   SpecialTypeKind Kind;
 
   unsigned Hash;
 
   SpecialTypeDescriptor() = delete;
   SpecialTypeDescriptor(SpecialTypeKind K) : Kind(K) { Hash = Kind; }
DeviceEventTypeDescriptor()
       : SpecialTypeDescriptor(SpecialTypeKind::STK_DeviceEvent) {
     Hash = Kind;
   }
 
   static bool classof(const SpecialTypeDescriptor *TD) {
     return TD->Kind == SpecialTypeKind::STK_DeviceEvent;
   }
 };
+
+struct PointerTypeDescriptor : public SpecialTypeDescriptor {
+  const Type *ElementType;
+  unsigned AddressSpace;
+
+  PointerTypeDescriptor() = delete;
+  PointerTypeDescriptor(const Type *ElementType, unsigned AddressSpace)
+      : SpecialTypeDescriptor(SpecialTypeKind::STK_Pointer),
+        ElementType(ElementType), AddressSpace(AddressSpace) {
+    Hash = (DenseMapInfo<Type *>().getHashValue(ElementType) & 0xffff) ^
+           ((AddressSpace << 8) | Kind);
+  }
+
+  static bool classof(const SpecialTypeDescriptor *TD) {
+    return TD->Kind == SpecialTypeKind::STK_Pointer;
+  }
+};
 } // namespace SPIRV
 
 template <> struct DenseMapInfo<SPIRV::SpecialTypeDescriptor> {
   static inline SPIRV::SpecialTypeDescriptor getEmptyKey() {
     return SPIRV::SpecialTypeDescriptor(
         SPIRV::SpecialTypeDescriptor::STK_Empty);
   }
   static inline SPIRV::SpecialTypeDescriptor getTombstoneKey() {
class SPIRVGeneralDuplicatesTracker {
   template <typename T>
   void prebuildReg2Entry(SPIRVDuplicatesTracker<T> &DT,
                          SPIRVReg2EntryTy &Reg2Entry);
 
 public:
   void buildDepsGraph(std::vector<SPIRV::DTSortableEntry *> &Graph,
                       MachineModuleInfo *MMI);
 
-  void add(const Type *T, const MachineFunction *MF, Register R) {
-    TT.add(T, MF, R);
+  void add(const Type *Ty, const MachineFunction *MF, Register R) {
+    TT.add(Ty, MF, R);
+  }
+
+  void add(const Type *PointerElementType, unsigned AddressSpace,
+           const MachineFunction *MF, Register R) {
+    ST.add(SPIRV::PointerTypeDescriptor(PointerElementType, AddressSpace), MF,
+           R);
   }
 
   void add(const Constant *C, const MachineFunction *MF, Register R) {
     CT.add(C, MF, R);
   }
 
   void add(const GlobalVariable *GV, const MachineFunction *MF, Register R) {
     GT.add(GV, MF, R);
   }
 
   void add(const Function *F, const MachineFunction *MF, Register R) {
     FT.add(F, MF, R);
   }
 
   void add(const Argument *Arg, const MachineFunction *MF, Register R) {
     AT.add(Arg, MF, R);
   }
 
   void add(const SPIRV::SpecialTypeDescriptor &TD, const MachineFunction *MF,
            Register R) {
     ST.add(TD, MF, R);
   }
 
-  Register find(const Type *T, const MachineFunction *MF) {
-    return TT.find(const_cast<Type *>(T), MF);
+  Register find(const Type *Ty, const MachineFunction *MF) {
+    return TT.find(const_cast<Type *>(Ty), MF);
+  }
+
+  Register find(const Type *PointerElementType, unsigned AddressSpace,
+                const MachineFunction *MF) {
+    return ST.find(
+        SPIRV::PointerTypeDescriptor(PointerElementType, AddressSpace), MF);
   }
 
   Register find(const Constant *C, const MachineFunction *MF) {
     return CT.find(const_cast<Constant *>(C), MF);
   }
 
   Register find(const GlobalVariable *GV, const MachineFunction *MF) {
     return GT.find(const_cast<GlobalVariable *>(GV), MF);

llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp

CallInst *buildIntrWithMD(Intrinsic::ID IntrID, ArrayRef<Type *> Types,
                             Value *Arg, Value *Arg2) {
     ConstantAsMetadata *CM = ValueAsMetadata::getConstant(Arg);
     MDTuple *TyMD = MDNode::get(F->getContext(), CM);
     MetadataAsValue *VMD = MetadataAsValue::get(F->getContext(), TyMD);
     return IRB->CreateIntrinsic(IntrID, {Types}, {Arg2, VMD});
   }
   void replaceMemInstrUses(Instruction *Old, Instruction *New);
   void processInstrAfterVisit(Instruction *I);
+  void insertAssignPtrTypeIntrs(Instruction *I);
   void insertAssignTypeIntrs(Instruction *I);
   void processGlobalValue(GlobalVariable &GV);
 
 public:
   static char ID;
   SPIRVEmitIntrinsics() : FunctionPass(ID) {
     initializeSPIRVEmitIntrinsicsPass(*PassRegistry::getPassRegistry());
   }
auto *InitInst = IRB->CreateIntrinsic(Intrinsic::spv_init_global,
                                           {GV.getType(), Ty}, {&GV, Const});
     InitInst->setArgOperand(1, Init);
   }
   if ((!GV.hasInitializer() || isa<UndefValue>(GV.getInitializer())) &&
       GV.getNumUses() == 0)
     IRB->CreateIntrinsic(Intrinsic::spv_unref_global, GV.getType(), &GV);
 }
 
+void SPIRVEmitIntrinsics::insertAssignPtrTypeIntrs(Instruction *I) {
+  if (I->getType()->isVoidTy() || !requireAssignType(I))
+    return;
+
+  setInsertPointSkippingPhis(*IRB, I->getNextNode());
+  if (auto *AI = dyn_cast<AllocaInst>(I)) {
+    Constant *Const = Constant::getNullValue(AI->getAllocatedType());
+    buildIntrWithMD(Intrinsic::spv_assign_ptr_type, {I->getType()}, Const, I);
+  }
+}
+
 void SPIRVEmitIntrinsics::insertAssignTypeIntrs(Instruction *I) {
   Type *Ty = I->getType();
-  if (!Ty->isVoidTy() && requireAssignType(I)) {
+  if (!Ty->isVoidTy() && requireAssignType(I) &&
+      I->getOpcode() != Instruction::Alloca) {
     setInsertPointSkippingPhis(*IRB, I->getNextNode());
     Type *TypeToAssign = Ty;
     if (auto *II = dyn_cast<IntrinsicInst>(I)) {
       if (II->getIntrinsicID() == Intrinsic::spv_const_composite ||
           II->getIntrinsicID() == Intrinsic::spv_undef) {
         auto t = AggrConsts.find(II);
         assert(t != AggrConsts.end());
         TypeToAssign = t->second->getType();
for (auto &GV : Func.getParent()->globals())
     processGlobalValue(GV);
 
   preprocessUndefs();
   preprocessCompositeConstants();
   SmallVector<Instruction *> Worklist;
   for (auto &I : instructions(Func))
     Worklist.push_back(&I);
 
-  for (auto &I : Worklist)
+  for (auto &I : Worklist) {
+    insertAssignPtrTypeIntrs(I);
     insertAssignTypeIntrs(I);
+  }
 
   for (auto *I : Worklist) {
     TrackConstants = true;
     if (!I->getType()->isVoidTy() || isa<StoreInst>(I))
       IRB->SetInsertPoint(I->getNextNode());
     I = visit(*I);
     processInstrAfterVisit(I);
   }
   return true;
 }
 
 FunctionPass *llvm::createSPIRVEmitIntrinsicsPass(SPIRVTargetMachine *TM) {
   return new SPIRVEmitIntrinsics(TM);
 }

llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp

 SPIRVType *SPIRVGlobalRegistry::getOrCreateOpTypeFunctionWithArgs(
     const Type *Ty, SPIRVType *RetType,
     const SmallVectorImpl<SPIRVType *> &ArgTypes,
     MachineIRBuilder &MIRBuilder) {
   Register Reg = DT.find(Ty, &MIRBuilder.getMF());
   if (Reg.isValid())
     return getSPIRVTypeForVReg(Reg);
   SPIRVType *SpirvType = getOpTypeFunction(RetType, ArgTypes, MIRBuilder);
+  DT.add(Ty, CurMF, getSPIRVTypeID(SpirvType));
   return finishCreatingSPIRVType(Ty, SpirvType);
 }
 
 SPIRVType *SPIRVGlobalRegistry::findSPIRVType(
     const Type *Ty, MachineIRBuilder &MIRBuilder,
     SPIRV::AccessQualifier::AccessQualifier AccQual, bool EmitIR) {
   Register Reg = DT.find(Ty, &MIRBuilder.getMF());
   if (Reg.isValid())
SPIRVType *SPIRVGlobalRegistry::restOfCreateSPIRVType(
   SPIRVType *SpirvType = createSPIRVType(Ty, MIRBuilder, AccessQual, EmitIR);
   TypesInProcessing.erase(Ty);
   VRegToTypeMap[&MIRBuilder.getMF()][getSPIRVTypeID(SpirvType)] = SpirvType;
   SPIRVToLLVMType[SpirvType] = Ty;
   Register Reg = DT.find(Ty, &MIRBuilder.getMF());
   // Do not add OpTypeForwardPointer to DT, a corresponding normal pointer type
   // will be added later. For special types it is already added to DT.
   if (SpirvType->getOpcode() != SPIRV::OpTypeForwardPointer && !Reg.isValid() &&
-      !isSpecialOpaqueType(Ty))
+      !isSpecialOpaqueType(Ty)) {
+    if (!Ty->isPointerTy())
       DT.add(Ty, &MIRBuilder.getMF(), getSPIRVTypeID(SpirvType));
+    else if (Ty->isOpaquePointerTy())
+      DT.add(Type::getInt8Ty(MIRBuilder.getMF().getFunction().getContext()),
+             Ty->getPointerAddressSpace(), &MIRBuilder.getMF(),
+             getSPIRVTypeID(SpirvType));
+    else
+      DT.add(Ty->getNonOpaquePointerElementType(), Ty->getPointerAddressSpace(),
+             &MIRBuilder.getMF(), getSPIRVTypeID(SpirvType));
+  }
 
   return SpirvType;
 }
 
 SPIRVType *SPIRVGlobalRegistry::getSPIRVTypeForVReg(Register VReg) const {
   auto t = VRegToTypeMap.find(CurMF);
   if (t != VRegToTypeMap.end()) {
     auto tt = t->second.find(VReg);
     if (tt != t->second.end())
       return tt->second;
   }
   return nullptr;
 }
 
 SPIRVType *SPIRVGlobalRegistry::getOrCreateSPIRVType(
     const Type *Ty, MachineIRBuilder &MIRBuilder,
     SPIRV::AccessQualifier::AccessQualifier AccessQual, bool EmitIR) {
-  Register Reg = DT.find(Ty, &MIRBuilder.getMF());
+  Register Reg;
+  if (!Ty->isPointerTy())
+    Reg = DT.find(Ty, &MIRBuilder.getMF());
+  else
+    Reg =
+        DT.find(Type::getInt8Ty(MIRBuilder.getMF().getFunction().getContext()),
+                Ty->getPointerAddressSpace(), &MIRBuilder.getMF());
+
   if (Reg.isValid() && !isSpecialOpaqueType(Ty))
     return getSPIRVTypeForVReg(Reg);
   TypesInProcessing.clear();
   SPIRVType *STy = restOfCreateSPIRVType(Ty, MIRBuilder, AccessQual, EmitIR);
   // Create normal pointer types for the corresponding OpTypeForwardPointers.
   for (auto &CU : ForwardPointerTypes) {
     const Type *Ty2 = CU.first;
     SPIRVType *STy2 = CU.second;

 }
 
 SPIRVType *SPIRVGlobalRegistry::getOrCreateOpTypeByOpcode(
     const Type *Ty, MachineIRBuilder &MIRBuilder, unsigned Opcode) {
   Register ResVReg = DT.find(Ty, &MIRBuilder.getMF());
   if (ResVReg.isValid())
     return MIRBuilder.getMF().getRegInfo().getUniqueVRegDef(ResVReg);
   ResVReg = createTypeVReg(MIRBuilder);
+  SPIRVType *SpirvTy = MIRBuilder.buildInstr(Opcode).addDef(ResVReg);
   DT.add(Ty, &MIRBuilder.getMF(), ResVReg);
-  return MIRBuilder.buildInstr(Opcode).addDef(ResVReg);
+  return SpirvTy;
 }
 
 const MachineInstr *
 SPIRVGlobalRegistry::checkSpecialInstr(const SPIRV::SpecialTypeDescriptor &TD,
                                        MachineIRBuilder &MIRBuilder) {
   Register Reg = DT.find(TD, &MIRBuilder.getMF());
   if (Reg.isValid())
     return MIRBuilder.getMF().getRegInfo().getUniqueVRegDef(Reg);
return getOrCreateSPIRVType(
       MIRBuilder);
 }
 
 SPIRVType *SPIRVGlobalRegistry::finishCreatingSPIRVType(const Type *LLVMTy,
                                                         SPIRVType *SpirvType) {
   assert(CurMF == SpirvType->getMF());
   VRegToTypeMap[CurMF][getSPIRVTypeID(SpirvType)] = SpirvType;
   SPIRVToLLVMType[SpirvType] = LLVMTy;
-  DT.add(LLVMTy, CurMF, getSPIRVTypeID(SpirvType));
   return SpirvType;
 }
 
 SPIRVType *SPIRVGlobalRegistry::getOrCreateSPIRVIntegerType(
     unsigned BitWidth, MachineInstr &I, const SPIRVInstrInfo &TII) {
   Type *LLVMTy = IntegerType::get(CurMF->getFunction().getContext(), BitWidth);
   Register Reg = DT.find(LLVMTy, CurMF);
   if (Reg.isValid())
     return getSPIRVTypeForVReg(Reg);
   MachineBasicBlock &BB = *I.getParent();
   auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpTypeInt))
                  .addDef(createTypeVReg(CurMF->getRegInfo()))
                  .addImm(BitWidth)
                  .addImm(0);
+  DT.add(LLVMTy, CurMF, getSPIRVTypeID(MIB));
   return finishCreatingSPIRVType(LLVMTy, MIB);
 }
 
 SPIRVType *
 SPIRVGlobalRegistry::getOrCreateSPIRVBoolType(MachineIRBuilder &MIRBuilder) {
   return getOrCreateSPIRVType(
       IntegerType::get(MIRBuilder.getMF().getFunction().getContext(), 1),
       MIRBuilder);
 }
 
 SPIRVType *
 SPIRVGlobalRegistry::getOrCreateSPIRVBoolType(MachineInstr &I,
                                               const SPIRVInstrInfo &TII) {
   Type *LLVMTy = IntegerType::get(CurMF->getFunction().getContext(), 1);
   Register Reg = DT.find(LLVMTy, CurMF);
   if (Reg.isValid())
     return getSPIRVTypeForVReg(Reg);
   MachineBasicBlock &BB = *I.getParent();
   auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpTypeBool))
                  .addDef(createTypeVReg(CurMF->getRegInfo()));
+  DT.add(LLVMTy, CurMF, getSPIRVTypeID(MIB));
   return finishCreatingSPIRVType(LLVMTy, MIB);
 }
 
 SPIRVType *SPIRVGlobalRegistry::getOrCreateSPIRVVectorType(
     SPIRVType *BaseType, unsigned NumElements, MachineIRBuilder &MIRBuilder) {
   return getOrCreateSPIRVType(
       FixedVectorType::get(const_cast<Type *>(getTypeForSPIRVType(BaseType)),
                            NumElements),
       MIRBuilder);
 }
 
 SPIRVType *SPIRVGlobalRegistry::getOrCreateSPIRVVectorType(
     SPIRVType *BaseType, unsigned NumElements, MachineInstr &I,
     const SPIRVInstrInfo &TII) {
   Type *LLVMTy = FixedVectorType::get(
       const_cast<Type *>(getTypeForSPIRVType(BaseType)), NumElements);
   Register Reg = DT.find(LLVMTy, CurMF);
   if (Reg.isValid())
     return getSPIRVTypeForVReg(Reg);
   MachineBasicBlock &BB = *I.getParent();
   auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpTypeVector))
                  .addDef(createTypeVReg(CurMF->getRegInfo()))
                  .addUse(getSPIRVTypeID(BaseType))
                  .addImm(NumElements);
+  DT.add(LLVMTy, CurMF, getSPIRVTypeID(MIB));
   return finishCreatingSPIRVType(LLVMTy, MIB);
 }
 
 SPIRVType *SPIRVGlobalRegistry::getOrCreateSPIRVArrayType(
     SPIRVType *BaseType, unsigned NumElements, MachineInstr &I,
     const SPIRVInstrInfo &TII) {
   Type *LLVMTy = ArrayType::get(
       const_cast<Type *>(getTypeForSPIRVType(BaseType)), NumElements);
   Register Reg = DT.find(LLVMTy, CurMF);
   if (Reg.isValid())
     return getSPIRVTypeForVReg(Reg);
   MachineBasicBlock &BB = *I.getParent();
   SPIRVType *SpirvType = getOrCreateSPIRVIntegerType(32, I, TII);
   Register Len = getOrCreateConstInt(NumElements, I, SpirvType, TII);
   auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpTypeArray))
                  .addDef(createTypeVReg(CurMF->getRegInfo()))
                  .addUse(getSPIRVTypeID(BaseType))
                  .addUse(Len);
+  DT.add(LLVMTy, CurMF, getSPIRVTypeID(MIB));
   return finishCreatingSPIRVType(LLVMTy, MIB);
 }
 
 SPIRVType *SPIRVGlobalRegistry::getOrCreateSPIRVPointerType(
     SPIRVType *BaseType, MachineIRBuilder &MIRBuilder,
-    SPIRV::StorageClass::StorageClass SClass) {
-  return getOrCreateSPIRVType(
-      PointerType::get(const_cast<Type *>(getTypeForSPIRVType(BaseType)),
-                       storageClassToAddressSpace(SClass)),
-      MIRBuilder);
+    SPIRV::StorageClass::StorageClass SC) {
+  const Type *PointerElementType = getTypeForSPIRVType(BaseType);
+  unsigned AddressSpace = storageClassToAddressSpace(SC);
+  Type *LLVMTy =
+      PointerType::get(const_cast<Type *>(PointerElementType), AddressSpace);
+  Register Reg = DT.find(PointerElementType, AddressSpace, CurMF);
+  if (Reg.isValid())
+    return getSPIRVTypeForVReg(Reg);
+  auto MIB = BuildMI(MIRBuilder.getMBB(), MIRBuilder.getInsertPt(),
+                     MIRBuilder.getDebugLoc(),
+                     MIRBuilder.getTII().get(SPIRV::OpTypePointer))
+                 .addDef(createTypeVReg(CurMF->getRegInfo()))
+                 .addImm(static_cast<uint32_t>(SC))
+                 .addUse(getSPIRVTypeID(BaseType));
+  DT.add(PointerElementType, AddressSpace, CurMF, getSPIRVTypeID(MIB));
+  return finishCreatingSPIRVType(LLVMTy, MIB);
 }
 
 SPIRVType *SPIRVGlobalRegistry::getOrCreateSPIRVPointerType(
     SPIRVType *BaseType, MachineInstr &I, const SPIRVInstrInfo &TII,
     SPIRV::StorageClass::StorageClass SC) {
+  const Type *PointerElementType = getTypeForSPIRVType(BaseType);
+  unsigned AddressSpace = storageClassToAddressSpace(SC);
   Type *LLVMTy =
-      PointerType::get(const_cast<Type *>(getTypeForSPIRVType(BaseType)),
-                       storageClassToAddressSpace(SC));
-  Register Reg = DT.find(LLVMTy, CurMF);
+      PointerType::get(const_cast<Type *>(PointerElementType), AddressSpace);
+  Register Reg = DT.find(PointerElementType, AddressSpace, CurMF);
   if (Reg.isValid())
     return getSPIRVTypeForVReg(Reg);
   MachineBasicBlock &BB = *I.getParent();
   auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpTypePointer))
                  .addDef(createTypeVReg(CurMF->getRegInfo()))
                  .addImm(static_cast<uint32_t>(SC))
                  .addUse(getSPIRVTypeID(BaseType));
+  DT.add(PointerElementType, AddressSpace, CurMF, getSPIRVTypeID(MIB));
   return finishCreatingSPIRVType(LLVMTy, MIB);
 }
 
 Register SPIRVGlobalRegistry::getOrCreateUndef(MachineInstr &I,
                                                SPIRVType *SpvType,
                                                const SPIRVInstrInfo &TII) {
   assert(SpvType);
   const Type *LLVMTy = getTypeForSPIRVType(SpvType);

llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp

bool SPIRVInstructionSelector::selectPhi(Register ResVReg,
   return MIB.constrainAllUses(TII, TRI, RBI);
 }
 
 bool SPIRVInstructionSelector::selectGlobalValue(
     Register ResVReg, MachineInstr &I, const MachineInstr *Init) const {
   // FIXME: don't use MachineIRBuilder here, replace it with BuildMI.
   MachineIRBuilder MIRBuilder(I);
   const GlobalValue *GV = I.getOperand(1).getGlobal();
-  SPIRVType *ResType = GR.getOrCreateSPIRVType(
-      GV->getType(), MIRBuilder, SPIRV::AccessQualifier::ReadWrite, false);
-
+  Type *GVType = GV->getValueType();
+  SPIRVType *PointerBaseType;
+  if (GVType->isArrayTy()) {
+    SPIRVType *ArrayElementType =
+        GR.getOrCreateSPIRVType(GVType->getArrayElementType(), MIRBuilder,
+                                SPIRV::AccessQualifier::ReadWrite, false);
+    PointerBaseType = GR.getOrCreateSPIRVArrayType(
+        ArrayElementType, GVType->getArrayNumElements(), I, TII);
+  } else {
+    PointerBaseType = GR.getOrCreateSPIRVType(
+        GVType, MIRBuilder, SPIRV::AccessQualifier::ReadWrite, false);
+  }
+  SPIRVType *ResType = GR.getOrCreateSPIRVPointerType(
+      PointerBaseType, I, TII,
+      addressSpaceToStorageClass(GV->getAddressSpace()));
   std::string GlobalIdent = GV->getGlobalIdentifier();
   // We have functions as operands in tests with blocks of instruction e.g. in
   // transcoding/global_block.ll. These operands are not used and should be
   // substituted by zero constants. Their type is expected to be always
   // OpTypePointer Function %uchar.
   if (isa<Function>(GV)) {
     const Constant *ConstVal = GV;
     MachineBasicBlock &BB = *I.getParent();
     Register NewReg = GR.find(ConstVal, GR.CurMF);
     if (!NewReg.isValid()) {
-      SPIRVType *SpvBaseTy = GR.getOrCreateSPIRVIntegerType(8, I, TII);
-      ResType = GR.getOrCreateSPIRVPointerType(SpvBaseTy, I, TII);
       Register NewReg = ResVReg;
       GR.add(ConstVal, GR.CurMF, NewReg);
       return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantNull))
           .addDef(NewReg)
           .addUse(GR.getSPIRVTypeID(ResType))
           .constrainAllUses(TII, TRI, RBI);
     }
     assert(NewReg != ResVReg);

llvm/lib/Target/SPIRV/SPIRVPreLegalizer.cpp

for (MachineBasicBlock *MBB : post_order(&MF)) {
     if (MBB->empty())
       continue;
 
     bool ReachedBegin = false;
     for (auto MII = std::prev(MBB->end()), Begin = MBB->begin();
          !ReachedBegin;) {
       MachineInstr &MI = *MII;
 
-      if (isSpvIntrinsic(MI, Intrinsic::spv_assign_type)) {
+      if (isSpvIntrinsic(MI, Intrinsic::spv_assign_ptr_type)) {
+        Register Reg = MI.getOperand(1).getReg();
+        MIB.setInsertPt(*MI.getParent(), MI.getIterator());
+        SPIRVType *BaseTy = GR->getOrCreateSPIRVType(
+            getMDOperandAsType(MI.getOperand(2).getMetadata(), 0), MIB);
+        SPIRVType *AssignedPtrType = GR->getOrCreateSPIRVPointerType(
+            BaseTy, MI, *MF.getSubtarget<SPIRVSubtarget>().getInstrInfo());
+        MachineInstr *Def = MRI.getVRegDef(Reg);
+        assert(Def && "Expecting an instruction that defines the register");
+        insertAssignInstr(Reg, nullptr, AssignedPtrType, GR, MIB,
+                          MF.getRegInfo());
+        ToErase.push_back(&MI);
+      } else if (isSpvIntrinsic(MI, Intrinsic::spv_assign_type)) {
         Register Reg = MI.getOperand(1).getReg();
         Type *Ty = getMDOperandAsType(MI.getOperand(2).getMetadata(), 0);
         MachineInstr *Def = MRI.getVRegDef(Reg);
         assert(Def && "Expecting an instruction that defines the register");
         // G_GLOBAL_VALUE already has type info.
         if (Def->getOpcode() != TargetOpcode::G_GLOBAL_VALUE)
           insertAssignInstr(Reg, Ty, nullptr, GR, MIB, MF.getRegInfo());
         ToErase.push_back(&MI);

llvm/test/CodeGen/SPIRV/EnqueueEmptyKernel.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
 
 ;; This test checks that Invoke parameter of OpEnueueKernel instruction meet the
 ;; following specification requirements in case of enqueueing empty block:
 ;; "Invoke must be an OpFunction whose OpTypeFunction operand has:
 ;; - Result Type must be OpTypeVoid.
 ;; - The first parameter must have a type of OpTypePointer to an 8-bit OpTypeInt.
 ;; - An optional list of parameters, each of which must have a type of OpTypePointer to the Workgroup Storage Class.
 ;; ... "

llvm/test/CodeGen/SPIRV/OpenCL/basic/vstore_private.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
 
 ; CHECK: %[[#i16_ty:]] = OpTypeInt 16 0
 ; CHECK: %[[#v4xi16_ty:]] = OpTypeVector %[[#i16_ty]] 4
 ; CHECK: %[[#pv4xi16_ty:]] = OpTypePointer Function %[[#v4xi16_ty]]
 ; CHECK: %[[#i16_const0:]] = OpConstant %[[#i16_ty]] 0
 ; CHECK: %[[#i16_undef:]] = OpUndef %[[#i16_ty]]
 ; CHECK: %[[#comp_const:]] = OpConstantComposite %[[#v4xi16_ty]] %[[#i16_const0]] %[[#i16_const0]] %[[#i16_const0]] %[[#i16_undef]]
 

llvm/test/CodeGen/SPIRV/OpenCL/get_global_id.ll

-; RUN: llc -O0 -mtriple=spirv64-unknown-unknown -opaque-pointers=0 %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
 
 ;; The set of valid inputs for get_global_id depends on the runtime NDRange,
 ;; but inputs outside of [0, 2] always return 0.
 ;; Here we assume Itanium mangling for function name.
 declare i64 @_Z13get_global_idj(i32)
 
 define i64 @foo(i32 %dim) {
   %x = call i64 @_Z13get_global_idj(i32 0)

llvm/test/CodeGen/SPIRV/OpenCL/get_global_offset.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
 
 ; CHECK: OpEntryPoint Kernel %[[#test_func:]] "test"
 ; CHECK: OpName %[[#outOffsets:]] "outOffsets"
 ; CHECK: OpName %[[#test_func]] "test"
 ; CHECK: OpName %[[#f2_decl:]] "BuiltInGlobalOffset"
 ; CHECK: OpDecorate %[[#f2_decl]] LinkageAttributes "BuiltInGlobalOffset" Import
-; CHECK: %[[#int_ty:]] = OpTypeInt 32 0
-; CHECK: %[[#iptr_ty:]] = OpTypePointer CrossWorkgroup  %[[#int_ty]]
+; CHECK: %[[#int_ty:]] = OpTypeInt 8 0
 ; CHECK: %[[#void_ty:]] = OpTypeVoid
+; CHECK: %[[#iptr_ty:]] = OpTypePointer CrossWorkgroup  %[[#int_ty]]
 ; CHECK: %[[#func_ty:]] = OpTypeFunction %[[#void_ty]] %[[#iptr_ty]]
 ; CHECK: %[[#int64_ty:]] = OpTypeInt 64 0
 ; CHECK: %[[#vec_ty:]] = OpTypeVector %[[#int64_ty]] 3
 ; CHECK: %[[#func2_ty:]] = OpTypeFunction %[[#vec_ty]]
 ;; TODO: add 64-bit constant defs
 ; CHECK: %[[#f2_decl]] = OpFunction %[[#vec_ty]] Pure %[[#func2_ty]]
 ; CHECK: OpFunctionEnd
 ;; Check that the function register name does not match other registers

llvm/test/CodeGen/SPIRV/OpenCL/progvar_prog_scope_init.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
 
 ; CHECK:     OpEntryPoint Kernel %[[#f1:]] "writer"
 ; CHECK:     OpEntryPoint Kernel %[[#f2:]] "reader"
 ; CHECK-DAG: OpName %[[#a_var:]] "a_var"
 ; CHECK-DAG: OpName %[[#p_var:]] "p_var"
 ; CHECK-DAG: %[[#uchar:]] = OpTypeInt 8 0
 ; CHECK-DAG: %[[#pt1:]] = OpTypePointer CrossWorkgroup %[[#uchar]]
 ; CHECK-DAG: %[[#arr2:]] = OpTypeArray %[[#uchar]]

llvm/test/CodeGen/SPIRV/OpenCL/wait_group_events.ll

 ; 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)
+define dso_local spir_kernel void @test_fn(ptr addrspace(1) %src) {
+  %src.addr = alloca ptr addrspace(1), align 8
+  store ptr addrspace(1) %src, ptr %src.addr, align 8
+  call spir_func void @_Z17wait_group_eventsiPU3AS49ocl_event(i32 0, ptr addrspace(4) null)
   ret void
 }
 
-declare spir_func void @_Z17wait_group_eventsiPU3AS49ocl_event(i32 noundef, %opencl.event_t* addrspace(4)* noundef)
+declare spir_func void @_Z17wait_group_eventsiPU3AS49ocl_event(i32, ptr addrspace(4))

llvm/test/CodeGen/SPIRV/SampledImageRetType.ll

-; RUN: llc -O0 -opaque-pointers=0 -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
-
-@__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
-}

llvm/test/CodeGen/SPIRV/atomicrmw.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
 
 ; CHECK:     %[[#Int:]] = OpTypeInt 32 0
 ; CHECK-DAG: %[[#Scope_Device:]] = OpConstant %[[#Int]] 1 {{$}}
 ; CHECK-DAG: %[[#MemSem_Relaxed:]] = OpConstant %[[#Int]] 0
 ; CHECK-DAG: %[[#MemSem_Acquire:]] = OpConstant %[[#Int]] 2
 ; CHECK-DAG: %[[#MemSem_Release:]] = OpConstant %[[#Int]] 4 {{$}}
 ; CHECK-DAG: %[[#MemSem_AcquireRelease:]] = OpConstant %[[#Int]] 8
 ; CHECK-DAG: %[[#MemSem_SequentiallyConsistent:]] = OpConstant %[[#Int]] 16

llvm/test/CodeGen/SPIRV/capabilities/capability-Float64-ImageBasic.ll

 ;; OpenCL C source
 ;; -----------------------------------------------
 ;; double d = 1.0;
 ;; kernel void test(read_only image2d_t img) {}
 ;; -----------------------------------------------
 
-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
-
-%opencl.image2d_t = type opaque
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
 
 @d = addrspace(1) global double 1.000000e+00, align 8
 
-define spir_kernel void @test(%opencl.image2d_t addrspace(1)* nocapture %img) {
+define spir_kernel void @test(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %img) {
 entry:
   ret void
 }
 
 ; CHECK-SPIRV-DAG: OpCapability Float64
 ; CHECK-SPIRV-DAG: OpCapability ImageBasic

llvm/test/CodeGen/SPIRV/capabilities/no_capability_shader.ll

 ;; __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:
+define spir_kernel void @sample_test(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %src, target("spirv.Image", void, 5, 0, 0, 0, 0, 0, 0) %buf) {
   ret void
 }

llvm/test/CodeGen/SPIRV/constants/global-constants.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
 
 @global   = addrspace(1) constant i32 1 ; OpenCL global memory
 @constant = addrspace(2) constant i32 2 ; OpenCL constant memory
 @local    = addrspace(3) constant i32 3 ; OpenCL local memory
 
 define i32 @getGlobal1() {
   %g = load i32, i32 addrspace(1)* @global
   ret i32 %g

llvm/test/CodeGen/SPIRV/constants/local-null-constants.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
 
 ;; OpenCL global memory
-define i32 addrspace(1)* @getConstant1() {
-  ret i32 addrspace(1)* null
+define ptr addrspace(1) @getConstant1() {
+  ret ptr addrspace(1) null
 }
 
 ;; OpenCL constant memory
-define i32 addrspace(2)* @getConstant2() {
-  ret i32 addrspace(2)* null
+define ptr addrspace(2) @getConstant2() {
+  ret ptr addrspace(2) null
 }
 
 ;; OpenCL local memory
-define i32 addrspace(3)* @getConstant3() {
-  ret i32 addrspace(3)* null
+define ptr addrspace(3) @getConstant3() {
+  ret ptr addrspace(3) null
 }
 
-; CHECK:     [[INT:%.+]] = OpTypeInt 32
+; CHECK:     [[INT:%.+]] = OpTypeInt 8
 
 ; CHECK-DAG: [[PTR_AS1:%.+]] = OpTypePointer CrossWorkgroup [[INT]]
 ; CHECK-DAG: OpConstantNull [[PTR_AS1]]
 
 ; CHECK-DAG: [[PTR_AS2:%.+]] = OpTypePointer UniformConstant [[INT]]
 ; CHECK-DAG: OpConstantNull [[PTR_AS2]]
 
 ; CHECK-DAG: [[PTR_AS3:%.+]] = OpTypePointer Workgroup [[INT]]
 ; CHECK-DAG: OpConstantNull [[PTR_AS3]]

llvm/test/CodeGen/SPIRV/function/alloca-load-store.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
+
+target triple = "spirv32-unknown-unknown"
 
 ; CHECK-DAG: OpName %[[#BAR:]] "bar"
 ; CHECK-DAG: OpName %[[#FOO:]] "foo"
 ; CHECK-DAG: OpName %[[#GOO:]] "goo"
 
-; CHECK:     %[[#INT:]] = OpTypeInt 32
+; CHECK-DAG: %[[#CHAR:]] = OpTypeInt 8
+; CHECK-DAG: %[[#INT:]] = OpTypeInt 32
 ; CHECK-DAG: %[[#STACK_PTR:]] = OpTypePointer Function %[[#INT]]
-; CHECK-DAG: %[[#GLOBAL_PTR:]] = OpTypePointer CrossWorkgroup %[[#INT]]
+; CHECK-DAG: %[[#GLOBAL_PTR:]] = OpTypePointer CrossWorkgroup %[[#CHAR]]
 ; CHECK-DAG: %[[#FN1:]] = OpTypeFunction %[[#INT]] %[[#INT]]
 ; CHECK-DAG: %[[#FN2:]] = OpTypeFunction %[[#INT]] %[[#INT]] %[[#GLOBAL_PTR]]
 
 define i32 @bar(i32 %a) {
   %p = alloca i32
   store i32 %a, i32* %p
   %b = load i32, i32* %p
   ret i32 %b

llvm/test/CodeGen/SPIRV/function/mangled-function.ll

 ;; 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)
+define spir_func void @bar(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %srcImage) local_unnamed_addr {
+  tail call spir_func void @_Z3foo14ocl_image2d_ro(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %srcImage)
   ret void
 }
 
-declare spir_func void @_Z3foo14ocl_image2d_ro(%opencl.image2d_ro_t addrspace(1)*) local_unnamed_addr
+declare spir_func void @_Z3foo14ocl_image2d_ro(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0)) local_unnamed_addr

llvm/test/CodeGen/SPIRV/half_extension.ll

 ;; #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 ;; half test()
 ;; {
 ;;   half x = 0.1f;
 ;;   x += 2.0f;
 ;;   half y = x + x;
 ;;   return y;
 ;; }
 
-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
 
 ; CHECK-SPIRV-DAG: OpCapability Float16Buffer
 ; CHECK-SPIRV-DAG: OpCapability Float16
 
 define spir_func half @test() {
 entry:
   %x = alloca half, align 2
   %y = alloca half, align 2

llvm/test/CodeGen/SPIRV/half_no_extension.ll

 ;; __kernel void test( __global float4 *p, __global half *f )
 ;; {
 ;;   __private float4 data;
 ;;   data = p[0];
 ;;   vstorea_half4_rtp( data, 0, f );
 ;; }
 
-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
 
 ; CHECK-SPIRV:     OpCapability Float16Buffer
 ; CHECK-SPIRV-NOT: OpCapability Float16
 
 define spir_kernel void @test(<4 x float> addrspace(1)* %p, half addrspace(1)* %f) {
 entry:
   %p.addr = alloca <4 x float> addrspace(1)*, align 8
   %f.addr = alloca half addrspace(1)*, align 8

llvm/test/CodeGen/SPIRV/image_decl_func_arg.ll

-; RUN: llc -O0 -opaque-pointers=0 -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
-}

llvm/test/CodeGen/SPIRV/image_dim.ll

-; RUN: llc -O0 -opaque-pointers=0 -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
-}

llvm/test/CodeGen/SPIRV/image_store.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
-
-;; Image types may be represented in two ways while translating to SPIR-V:
-;; - OpenCL form based on pointers-to-opaque-structs, e.g. '%opencl.image2d_ro_t',
-;; - SPIR-V form based on TargetExtType, e.g. 'target("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
-
-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 target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0), align 8
-  store %opencl.image2d_ro_t addrspace(1)* %srcimg, %opencl.image2d_ro_t addrspace(1)** %srcimg.addr, align 8
-  ret void
-}

llvm/test/CodeGen/SPIRV/link-attribute.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
-
-%opencl.image2d_t = type opaque
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
 
 ; CHECK: OpDecorate %[[#ID:]] LinkageAttributes "imageSampler" Export
 ; CHECK: %[[#ID]] = OpVariable %[[#]] UniformConstant %[[#]]
 
 @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) {
+define spir_kernel void @sample_kernel(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %input, ptr addrspace(1) nocapture %xOffsets, ptr addrspace(1) nocapture %yOffsets, ptr 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)
+  %5 = tail call spir_func i32 @_Z15get_image_width11ocl_image2d(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %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)
+  %12 = tail call spir_func <4 x float> @_Z11read_imagef11ocl_image2d11ocl_samplerDv2_f(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %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
+  %14 = getelementptr inbounds <4 x float>, ptr addrspace(1) %results, i64 %13
+  store <4 x float> %12, ptr 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 i32 @_Z15get_image_width11ocl_image2d(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0))
 
-declare spir_func <4 x float> @_Z11read_imagef11ocl_image2d11ocl_samplerDv2_f(%opencl.image2d_t addrspace(1)*, i32, <2 x float>)
+declare spir_func <4 x float> @_Z11read_imagef11ocl_image2d11ocl_samplerDv2_f(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0), i32, <2 x float>)

llvm/test/CodeGen/SPIRV/linked-list.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
-
-%struct.Node = type { %struct.Node.0 addrspace(1)* }
-; CHECK: %[[#]] = OpTypeOpaque "struct.Node.0"
-%struct.Node.0 = type opaque
-
-define spir_kernel void @create_linked_lists(%struct.Node addrspace(1)* nocapture %pNodes, i32 addrspace(1)* nocapture %allocation_index, i32 %list_length) {
-entry:
-  ret void
-}

llvm/test/CodeGen/SPIRV/literal-struct.ll

-;; This test checks that the backend doesn't crash if the module has literal
-;; structs, i.e. structs whose type has no name. Typicaly clang generate such
-;; structs if the kernel contains OpenCL 2.0 blocks. The IR was produced with
-;; the following command:
-;; clang -cc1 -triple spir -cl-std=cl2.0 -O0 literal-struct.cl -emit-llvm -o test/literal-struct.ll
-
-;; literal-struct.cl:
-;; void foo()
-;; {
-;;   void (^myBlock)(void) = ^{};
-;;   myBlock();
-;; }
-
-; RUN: llc -opaque-pointers=0 -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
-
-; CHECK: OpName %[[#StructType0:]] "struct.__opencl_block_literal_generic"
-; CHECK: %[[#Int8:]] = OpTypeInt 8 0
-; CHECK: %[[#Int8Ptr:]] = OpTypePointer Generic %[[#Int8]]
-; CHECK: %[[#Int:]] = OpTypeInt 32 0
-; CHECK: %[[#StructType0:]] = OpTypeStruct %[[#Int]] %[[#Int]] %[[#Int8Ptr]]
-; CHECK: %[[#StructType:]] = OpTypeStruct %[[#Int]] %[[#Int]] %[[#Int8Ptr]]
-
-%struct.__opencl_block_literal_generic = type { i32, i32, i8 addrspace(4)* }
-
-@__block_literal_global = internal addrspace(1) constant { i32, i32, i8 addrspace(4)* } { i32 12, i32 4, i8 addrspace(4)* addrspacecast (i8* bitcast (void (i8 addrspace(4)*)* @__foo_block_invoke to i8*) to i8 addrspace(4)*) }, align 4
-; CHECK: OpConstantComposite %[[#StructType]]
-
-@__block_literal_global.1 = internal addrspace(1) constant { i32, i32, i8 addrspace(4)* } zeroinitializer, align 4
-; CHECK: OpConstantNull %[[#StructType]]
-
-define spir_func void @foo() {
-entry:
-  %myBlock = alloca %struct.__opencl_block_literal_generic addrspace(4)*, align 4
-  store %struct.__opencl_block_literal_generic addrspace(4)* addrspacecast (%struct.__opencl_block_literal_generic addrspace(1)* bitcast ({ i32, i32, i8 addrspace(4)* } addrspace(1)* @__block_literal_global to %struct.__opencl_block_literal_generic addrspace(1)*) to %struct.__opencl_block_literal_generic addrspace(4)*), %struct.__opencl_block_literal_generic addrspace(4)** %myBlock, align 4
-  call spir_func void @__foo_block_invoke(i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast ({ i32, i32, i8 addrspace(4)* } addrspace(1)* @__block_literal_global to i8 addrspace(1)*) to i8 addrspace(4)*))
-  ret void
-}
-
-define internal spir_func void @__foo_block_invoke(i8 addrspace(4)* %.block_descriptor) {
-entry:
-  %.block_descriptor.addr = alloca i8 addrspace(4)*, align 4
-  %block.addr = alloca <{ i32, i32, i8 addrspace(4)* }> addrspace(4)*, align 4
-  store i8 addrspace(4)* %.block_descriptor, i8 addrspace(4)** %.block_descriptor.addr, align 4
-  %block = bitcast i8 addrspace(4)* %.block_descriptor to <{ i32, i32, i8 addrspace(4)* }> addrspace(4)*
-  store <{ i32, i32, i8 addrspace(4)* }> addrspace(4)* %block, <{ i32, i32, i8 addrspace(4)* }> addrspace(4)** %block.addr, align 4
-  ret void
-}

llvm/test/CodeGen/SPIRV/llvm-intrinsics/memset.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
 
-; CHECK-SPIRV: OpDecorate %[[#NonConstMemset:]] LinkageAttributes "spirv.llvm_memset_p3i8_i32"
-; CHECK-SPIRV: %[[#Int32:]] = OpTypeInt 32 0
-; CHECK-SPIRV: %[[#Int8:]] = OpTypeInt 8 0
-; CHECK-SPIRV: %[[#Int8Ptr:]] = OpTypePointer Generic %[[#Int8]]
-; CHECK-SPIRV: %[[#Lenmemset21:]] = OpConstant %[[#]] 4
-; CHECK-SPIRV: %[[#Int8x4:]] = OpTypeArray %[[#Int8]] %[[#Lenmemset21]]
-; CHECK-SPIRV: %[[#Int8PtrConst:]] = OpTypePointer UniformConstant %[[#Int8]]
-; CHECK-SPIRV: %[[#Lenmemset0:]] = OpConstant %[[#Int32]] 12
-; CHECK-SPIRV: %[[#Int8x12:]] = OpTypeArray %[[#Int8]] %[[#Lenmemset0]]
-; CHECK-SPIRV: %[[#Const21:]] = OpConstant %[[#]] 21
-; CHECK-SPIRV: %[[#False:]] = OpConstantFalse %[[#]]
-; CHECK-SPIRV: %[[#InitComp:]] = OpConstantComposite %[[#Int8x4]] %[[#Const21]] %[[#Const21]] %[[#Const21]] %[[#Const21]]
-; CHECK-SPIRV: %[[#Init:]] = OpConstantNull %[[#Int8x12]]
-; CHECK-SPIRV: %[[#ValComp:]] = OpVariable %[[#]] UniformConstant %[[#InitComp]]
-; CHECK-SPIRV: %[[#Val:]] = OpVariable %[[#]] UniformConstant %[[#Init]]
-
-; CHECK-SPIRV: %[[#Target:]] = OpBitcast %[[#Int8Ptr]] %[[#]]
-; CHECK-SPIRV: %[[#Source:]] = OpBitcast %[[#Int8PtrConst]] %[[#Val]]
-; CHECK-SPIRV: OpCopyMemorySized %[[#Target]] %[[#Source]] %[[#Lenmemset0]] Aligned 4
-
-; CHECK-SPIRV: %[[#SourceComp:]] = OpBitcast %[[#Int8PtrConst]] %[[#ValComp]]
-; CHECK-SPIRV: OpCopyMemorySized %[[#]] %[[#SourceComp]] %[[#Lenmemset21]] Aligned 4
-
-; CHECK-SPIRV: %[[#]] = OpFunctionCall %[[#]] %[[#NonConstMemset]] %[[#]] %[[#]] %[[#]] %[[#False]]
-
-; CHECK-SPIRV: %[[#NonConstMemset]] = OpFunction %[[#]]
-; CHECK-SPIRV: %[[#Dest:]] = OpFunctionParameter %[[#]]
-; CHECK-SPIRV: %[[#Value:]] = OpFunctionParameter %[[#]]
-; CHECK-SPIRV: %[[#Len:]] = OpFunctionParameter %[[#]]
-; CHECK-SPIRV: %[[#Volatile:]] = OpFunctionParameter %[[#]]
-
-; CHECK-SPIRV: %[[#Entry:]] = OpLabel
-; CHECK-SPIRV: %[[#IsZeroLen:]] = OpIEqual %[[#]] %[[#Zero:]] %[[#Len]]
-; CHECK-SPIRV: OpBranchConditional %[[#IsZeroLen]] %[[#End:]] %[[#WhileBody:]]
-
-; CHECK-SPIRV: %[[#WhileBody]] = OpLabel
-; CHECK-SPIRV: %[[#Offset:]] = OpPhi %[[#]] %[[#Zero]] %[[#Entry]] %[[#OffsetInc:]] %[[#WhileBody]]
-; CHECK-SPIRV: %[[#Ptr:]] = OpInBoundsPtrAccessChain %[[#]] %[[#Dest]] %[[#Offset]]
-; CHECK-SPIRV: OpStore %[[#Ptr]] %[[#Value]] Aligned 1
-; CHECK-SPIRV: %[[#OffsetInc]] = OpIAdd %[[#]] %[[#Offset]] %[[#One:]]
-; CHECK-SPIRV: %[[#NotEnd:]] = OpULessThan %[[#]] %[[#OffsetInc]] %[[#Len]]
-; CHECK-SPIRV: OpBranchConditional %[[#NotEnd]] %[[#WhileBody]] %[[#End]]
+; CHECK-DAG: OpDecorate %[[#Memset_p0i32:]] LinkageAttributes "spirv.llvm_memset_p0_i32" Export
+; CHECK-DAG: OpDecorate %[[#Memset_p3i32:]] LinkageAttributes "spirv.llvm_memset_p3_i32" Export
+; CHECK-DAG: OpDecorate %[[#Memset_p1i64:]] LinkageAttributes "spirv.llvm_memset_p1_i64" Export
+; CHECK-DAG: OpDecorate %[[#Memset_p1i64:]] LinkageAttributes "spirv.llvm_memset_p1_i64.volatile" Export
+
+; CHECK-DAG: %[[#Int8:]] = OpTypeInt 8 0
+; CHECK-DAG: %[[#Int32:]] = OpTypeInt 32 0
+; CHECK-DAG: %[[#Int64:]] = OpTypeInt 64 0
+; CHECK-DAG: %[[#Void:]] = OpTypeVoid
+; CHECK-DAG: %[[#Int8Ptr:]] = OpTypePointer Generic %[[#Int8]]
+
+; CHECK-DAG: %[[#Const4:]] = OpConstant %[[#Int32]] 4
+; CHECK: %[[#Int8x4:]] = OpTypeArray %[[#Int8]] %[[#Const4]]
+
+; CHECK-DAG: %[[#Const12:]] = OpConstant %[[#Int32]] 12
+; CHECK: %[[#Int8x12:]] = OpTypeArray %[[#Int8]] %[[#Const12]]
+
+; CHECK-DAG: %[[#Const21:]] = OpConstant %[[#Int8]] 21
+; CHECK-DAG: %[[#False:]] = OpConstantFalse %[[#]]
+; CHECK-DAG: %[[#ConstComp:]] = OpConstantComposite %[[#Int8x4]] %[[#Const21]] %[[#Const21]] %[[#Const21]] %[[#Const21]]
+; CHECK-DAG: %[[#ConstNull:]] = OpConstantNull %[[#Int8x12]]
+; CHECK: %[[#VarComp:]] = OpVariable %[[#]] UniformConstant %[[#ConstComp]]
+; CHECK: %[[#VarNull:]] = OpVariable %[[#]] UniformConstant %[[#ConstNull]]
+
+; CHECK-DAG: %[[#Int8PtrConst:]] = OpTypePointer UniformConstant %[[#Int8]]
+; CHECK: %[[#Target:]] = OpBitcast %[[#Int8Ptr]] %[[#]]
+; CHECK: %[[#Source:]] = OpBitcast %[[#Int8PtrConst]] %[[#VarNull]]
+; CHECK: OpCopyMemorySized %[[#Target]] %[[#Source]] %[[#Const12]] Aligned 4
+
+; CHECK: %[[#SourceComp:]] = OpBitcast %[[#Int8PtrConst]] %[[#VarComp]]
+; CHECK: OpCopyMemorySized %[[#]] %[[#SourceComp]] %[[#Const4]] Aligned 4
+
+; CHECK-SPIRV: %[[#]] = OpFunctionCall %[[#]] %[[#Memset_p0i32]] %[[#]] %[[#]] %[[#]] %[[#False]]
+
+; CHECK: %[[#Memset_p0i32]] = OpFunction %[[#]]
+; CHECK: %[[#Dest:]] = OpFunctionParameter %[[#]]
+; CHECK: %[[#Value:]] = OpFunctionParameter %[[#]]
+; CHECK: %[[#Len:]] = OpFunctionParameter %[[#]]
+; CHECK: %[[#Volatile:]] = OpFunctionParameter %[[#]]
+
+; CHECK: %[[#Entry:]] = OpLabel
+; CHECK: %[[#IsZeroLen:]] = OpIEqual %[[#]] %[[#Zero:]] %[[#Len]]
+; CHECK: OpBranchConditional %[[#IsZeroLen]] %[[#End:]] %[[#WhileBody:]]
+
+; CHECK: %[[#WhileBody]] = OpLabel
+; CHECK: %[[#Offset:]] = OpPhi %[[#]] %[[#Zero]] %[[#Entry]] %[[#OffsetInc:]] %[[#WhileBody]]
+; CHECK: %[[#Ptr:]] = OpInBoundsPtrAccessChain %[[#]] %[[#Dest]] %[[#Offset]]
+; CHECK: OpStore %[[#Ptr]] %[[#Value]] Aligned 1
+; CHECK: %[[#OffsetInc]] = OpIAdd %[[#]] %[[#Offset]] %[[#One:]]
+; CHECK: %[[#NotEnd:]] = OpULessThan %[[#]] %[[#OffsetInc]] %[[#Len]]
+; CHECK: OpBranchConditional %[[#NotEnd]] %[[#WhileBody]] %[[#End]]
 
-; CHECK-SPIRV: %[[#End]] = OpLabel
-; CHECK-SPIRV: OpReturn
+; CHECK: %[[#End]] = OpLabel
+; CHECK: OpReturn
 
-; CHECK-SPIRV: OpFunctionEnd
+; CHECK: OpFunctionEnd
 
 %struct.S1 = type { i32, i32, i32 }
 
 define spir_func void @_Z5foo11v(%struct.S1 addrspace(4)* noalias nocapture sret(%struct.S1 addrspace(4)*) %agg.result, i32 %s1, i64 %s2, i8 %v) {
   %x = alloca [4 x i8]
   %x.bc = bitcast [4 x i8]* %x to i8*
-  %1 = bitcast %struct.S1 addrspace(4)* %agg.result to i8 addrspace(4)*
-  tail call void @llvm.memset.p4i8.i32(i8 addrspace(4)* align 4 %1, i8 0, i32 12, i1 false)
+  %a = bitcast %struct.S1 addrspace(4)* %agg.result to i8 addrspace(4)*
+  tail call void @llvm.memset.p4i8.i32(i8 addrspace(4)* align 4 %a, i8 0, i32 12, i1 false)
   tail call void @llvm.memset.p0i8.i32(i8* align 4 %x.bc, i8 21, i32 4, i1 false)
 
   ;; non-const value
   tail call void @llvm.memset.p0i8.i32(i8* align 4 %x.bc, i8 %v, i32 3, i1 false)
 
   ;; non-const value and size
   tail call void @llvm.memset.p0i8.i32(i8*  align 4 %x.bc, i8 %v, i32 %s1, i1 false)
 
   ;; Address spaces, non-const value and size
-  %a = addrspacecast i8 addrspace(4)* %1 to i8 addrspace(3)*
-  tail call void @llvm.memset.p3i8.i32(i8 addrspace(3)* align 4 %a, i8 %v, i32 %s1, i1 false)
-  %b = addrspacecast i8 addrspace(4)* %1 to i8 addrspace(1)*
-  tail call void @llvm.memset.p1i8.i64(i8 addrspace(1)* align 4 %b, i8 %v, i64 %s2, i1 false)
+  %b = addrspacecast i8 addrspace(4)* %a to i8 addrspace(3)*
+  tail call void @llvm.memset.p3i8.i32(i8 addrspace(3)* align 4 %b, i8 %v, i32 %s1, i1 false)
+  %c = addrspacecast i8 addrspace(4)* %a to i8 addrspace(1)*
+  tail call void @llvm.memset.p1i8.i64(i8 addrspace(1)* align 4 %c, i8 %v, i64 %s2, i1 false)
 
   ;; Volatile
-  tail call void @llvm.memset.p1i8.i64(i8 addrspace(1)* align 4 %b, i8 %v, i64 %s2, i1 true)
+  tail call void @llvm.memset.p1i8.i64(i8 addrspace(1)* align 4 %c, i8 %v, i64 %s2, i1 true)
   ret void
 }
 
 declare void @llvm.memset.p4i8.i32(i8 addrspace(4)* nocapture, i8, i32, i1)
 
 declare void @llvm.memset.p0i8.i32(i8* nocapture, i8, i32, i1)
 
 declare void @llvm.memset.p3i8.i32(i8 addrspace(3)*, i8, i32, i1)
 
 declare void @llvm.memset.p1i8.i64(i8 addrspace(1)*, i8, i64, i1)

llvm/test/CodeGen/SPIRV/llvm-intrinsics/sqrt.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
+; 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]]
+; CHECK-DAG: %[[#ExtInstSetId:]] = OpExtInstImport "OpenCL.std"
+; CHECK-DAG: %[[#Float:]] = OpTypeFloat 32
+; CHECK-DAG: %[[#Double:]] = OpTypeFloat 64
+; CHECK-DAG: %[[#Double4:]] = OpTypeVector %[[#Double]] 4
+; CHECK-DAG: %[[#FloatArg:]] = OpConstant %[[#Float]]
+; CHECK-DAG: %[[#DoubleArg:]] = OpConstant %[[#Double]]
+; CHECK-DAG: %[[#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

llvm/test/CodeGen/SPIRV/opaque_pointers.ll

 ; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK
 
-; CHECK: %[[#Int8Ty:]] = OpTypeInt 8 0
-; CHECK: %[[#PtrTy:]] = OpTypePointer Function %[[#Int8Ty]]
-; CHECK: %[[#Int64Ty:]] = OpTypeInt 64 0
-; CHECK: %[[#FTy:]] = OpTypeFunction %[[#Int64Ty]] %[[#PtrTy]]
-; CHECK: %[[#Int32Ty:]] = OpTypeInt 32 0
-; CHECK: %[[#Const:]] = OpConstant %[[#Int32Ty]] 0
+; CHECK-DAG: %[[#Int8Ty:]] = OpTypeInt 8 0
+; CHECK-DAG: %[[#PtrTy:]] = OpTypePointer Function %[[#Int8Ty]]
+; CHECK-DAG: %[[#Int64Ty:]] = OpTypeInt 64 0
+; CHECK-DAG: %[[#FTy:]] = OpTypeFunction %[[#Int64Ty]] %[[#PtrTy]]
+; CHECK-DAG: %[[#Int32Ty:]] = OpTypeInt 32 0
+; CHECK-DAG: %[[#Const:]] = OpConstant %[[#Int32Ty]] 0
 ; CHECK: OpFunction %[[#Int64Ty]] None %[[#FTy]]
 ; CHECK: %[[#Parm:]] = OpFunctionParameter %[[#PtrTy]]
 ; CHECK: OpStore %[[#Parm]] %[[#Const]] Aligned 4
 ; CHECK: %[[#Res:]] = OpLoad %[[#Int64Ty]] %[[#Parm]] Aligned 8
 ; CHECK: OpReturnValue %[[#Res]]
 
 define i64 @test(ptr %p) {
   store i32 0, ptr %p
   %v = load i64, ptr %p
   ret i64 %v
 }

llvm/test/CodeGen/SPIRV/opencl.queue_t.ll

-; RUN: llc -O0 -opaque-pointers=0 -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
-}

llvm/test/CodeGen/SPIRV/pstruct.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
-
-%struct.ST = type { i32, i32, i32 }
-
-; CHECK-SPIRV: OpName %[[#struct:]] "struct.ST"
-; CHECK-SPIRV: %[[#int:]] = OpTypeInt 32 0
-; CHECK-SPIRV: %[[#intP:]] = OpTypePointer Function %[[#int]]
-; CHECK-SPIRV: %[[#struct]] = OpTypeStruct %[[#int]] %[[#int]] %[[#int]]
-; CHECK-SPIRV: %[[#structP:]] = OpTypePointer Function %[[#struct]]
-; CHECK-SPIRV: %[[#structPP:]] = OpTypePointer Function %[[#structP]]
-; CHECK-SPIRV: %[[#zero:]] = OpConstant %[[#int]] 0
-; CHECK-SPIRV: %[[#one:]] = OpConstant %[[#int]] 1
-; CHECK-SPIRV: %[[#two:]] = OpConstant %[[#int]] 2
-
-define dso_local spir_func i32 @cmp_func(i8* %p1, i8* %p2) {
-entry:
-  %retval = alloca i32, align 4
-  %p1.addr = alloca i8*, align 8
-  %p2.addr = alloca i8*, align 8
-; CHECK-SPIRV: %[[#s1:]] = OpVariable %[[#structPP]]
-; CHECK-SPIRV: %[[#s2:]] = OpVariable %[[#structPP]]
-  %s1 = alloca %struct.ST*, align 8
-  %s2 = alloca %struct.ST*, align 8
-  store i8* %p1, i8** %p1.addr, align 8
-  store i8* %p2, i8** %p2.addr, align 8
-  %0 = load i8*, i8** %p1.addr, align 8
-; CHECK-SPIRV: %[[#t1:]] = OpBitcast %[[#structP]]
-; CHECK-SPIRV: OpStore %[[#s1]] %[[#t1]]
-  %1 = bitcast i8* %0 to %struct.ST*
-  store %struct.ST* %1, %struct.ST** %s1, align 8
-  %2 = load i8*, i8** %p2.addr, align 8
-; CHECK-SPIRV: %[[#t2:]] = OpBitcast %[[#structP]]
-; CHECK-SPIRV: OpStore %[[#s2]] %[[#t2]]
-  %3 = bitcast i8* %2 to %struct.ST*
-  store %struct.ST* %3, %struct.ST** %s2, align 8
-; CHECK-SPIRV: %[[#t3:]] = OpLoad %[[#structP]] %[[#s1]]
-; CHECK-SPIRV: %[[#a1:]] = OpInBoundsPtrAccessChain %[[#intP]] %[[#t3]] %[[#zero]] %[[#zero]]
-; CHECK-SPIRV: %[[#]] = OpLoad %[[#int]] %[[#a1]]
-  %4 = load %struct.ST*, %struct.ST** %s1, align 8
-  %a = getelementptr inbounds %struct.ST, %struct.ST* %4, i32 0, i32 0
-  %5 = load i32, i32* %a, align 4
-; CHECK-SPIRV: %[[#t4:]] = OpLoad %[[#structP]] %[[#s2]]
-; CHECK-SPIRV: %[[#a2:]] = OpInBoundsPtrAccessChain %[[#intP]] %[[#t4]] %[[#zero]] %[[#zero]]
-; CHECK-SPIRV: %[[#]] = OpLoad %[[#int]] %[[#a2]]
-  %6 = load %struct.ST*, %struct.ST** %s2, align 8
-  %a1 = getelementptr inbounds %struct.ST, %struct.ST* %6, i32 0, i32 0
-  %7 = load i32, i32* %a1, align 4
-  %cmp = icmp ne i32 %5, %7
-  br i1 %cmp, label %if.then, label %if.end
-
-if.then:                                          ; preds = %entry
-; CHECK-SPIRV: %[[#t5:]] = OpLoad %[[#structP]] %[[#s1]]
-; CHECK-SPIRV: %[[#a_1:]] = OpInBoundsPtrAccessChain %[[#intP]] %[[#t5]] %[[#zero]] %[[#zero]]
-; CHECK-SPIRV: %[[#]] = OpLoad %[[#int]] %[[#a_1]]
-  %8 = load %struct.ST*, %struct.ST** %s1, align 8
-  %a2 = getelementptr inbounds %struct.ST, %struct.ST* %8, i32 0, i32 0
-  %9 = load i32, i32* %a2, align 4
-; CHECK-SPIRV: %[[#t6:]] = OpLoad %[[#structP]] %[[#s2]]
-; CHECK-SPIRV: %[[#a_2:]] = OpInBoundsPtrAccessChain %[[#intP]] %[[#t6]] %[[#zero]] %[[#zero]]
-; CHECK-SPIRV: %[[#]] = OpLoad %[[#int]] %[[#a_2]]
-  %10 = load %struct.ST*, %struct.ST** %s2, align 8
-  %a3 = getelementptr inbounds %struct.ST, %struct.ST* %10, i32 0, i32 0
-  %11 = load i32, i32* %a3, align 4
-  %sub = sub nsw i32 %9, %11
-  store i32 %sub, i32* %retval, align 4
-  br label %return
-
-if.end:                                           ; preds = %entry
-; CHECK-SPIRV: %[[#t7:]] = OpLoad %[[#structP]] %[[#s1]]
-; CHECK-SPIRV: %[[#b1:]] = OpInBoundsPtrAccessChain %[[#intP]] %[[#t7]] %[[#zero]] %[[#one]]
-; CHECK-SPIRV: %[[#]] = OpLoad %[[#int]] %[[#b1]]
-  %12 = load %struct.ST*, %struct.ST** %s1, align 8
-  %b = getelementptr inbounds %struct.ST, %struct.ST* %12, i32 0, i32 1
-  %13 = load i32, i32* %b, align 4
-; CHECK-SPIRV: %[[#t8:]] = OpLoad %[[#structP]] %[[#s2]]
-; CHECK-SPIRV: %[[#b2:]] = OpInBoundsPtrAccessChain %[[#intP]] %[[#t8]] %[[#zero]] %[[#one]]
-; CHECK-SPIRV: %[[#]] = OpLoad %[[#int]] %[[#b2]]
-  %14 = load %struct.ST*, %struct.ST** %s2, align 8
-  %b4 = getelementptr inbounds %struct.ST, %struct.ST* %14, i32 0, i32 1
-  %15 = load i32, i32* %b4, align 4
-  %cmp5 = icmp ne i32 %13, %15
-  br i1 %cmp5, label %if.then6, label %if.end10
-
-if.then6:                                         ; preds = %if.end
-; CHECK-SPIRV: %[[#t9:]] = OpLoad %[[#structP]] %[[#s1]]
-; CHECK-SPIRV: %[[#b_1:]] = OpInBoundsPtrAccessChain %[[#intP]] %[[#t9]] %[[#zero]] %[[#one]]
-; CHECK-SPIRV: %[[#]] = OpLoad %[[#int]] %[[#b_1]]
-  %16 = load %struct.ST*, %struct.ST** %s1, align 8
-  %b7 = getelementptr inbounds %struct.ST, %struct.ST* %16, i32 0, i32 1
-  %17 = load i32, i32* %b7, align 4
-; CHECK-SPIRV: %[[#t10:]] = OpLoad %[[#structP]] %[[#s2]]
-; CHECK-SPIRV: %[[#b_2:]] = OpInBoundsPtrAccessChain %[[#intP]] %[[#t10]] %[[#zero]] %[[#one]]
-; CHECK-SPIRV: %[[#]] = OpLoad %[[#int]] %[[#b_2]]
-  %18 = load %struct.ST*, %struct.ST** %s2, align 8
-  %b8 = getelementptr inbounds %struct.ST, %struct.ST* %18, i32 0, i32 1
-  %19 = load i32, i32* %b8, align 4
-  %sub9 = sub nsw i32 %17, %19
-  store i32 %sub9, i32* %retval, align 4
-  br label %return
-
-if.end10:                                         ; preds = %if.end
-; CHECK-SPIRV: %[[#t11:]] = OpLoad %[[#structP]] %[[#s1]]
-; CHECK-SPIRV: %[[#c1:]] = OpInBoundsPtrAccessChain %[[#intP]] %[[#t11]] %[[#zero]] %[[#two]]
-; CHECK-SPIRV: %[[#]] = OpLoad %[[#int]] %[[#c1]]
-  %20 = load %struct.ST*, %struct.ST** %s1, align 8
-  %c = getelementptr inbounds %struct.ST, %struct.ST* %20, i32 0, i32 2
-  %21 = load i32, i32* %c, align 4
-; CHECK-SPIRV: %[[#t12:]] = OpLoad %[[#structP]] %[[#s2]]
-; CHECK-SPIRV: %[[#c2:]] = OpInBoundsPtrAccessChain %[[#intP]] %[[#t12]] %[[#zero]] %[[#two]]
-; CHECK-SPIRV: %[[#]] = OpLoad %[[#int]] %[[#c2]]
-  %22 = load %struct.ST*, %struct.ST** %s2, align 8
-  %c11 = getelementptr inbounds %struct.ST, %struct.ST* %22, i32 0, i32 2
-  %23 = load i32, i32* %c11, align 4
-  %sub12 = sub nsw i32 %21, %23
-  store i32 %sub12, i32* %retval, align 4
-  br label %return
-
-return:                                           ; preds = %if.end10, %if.then6, %if.then
-  %24 = load i32, i32* %retval, align 4
-  ret i32 %24
-}

llvm/test/CodeGen/SPIRV/read_image.ll

-; RUN: llc -O0 -opaque-pointers=0 -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)

llvm/test/CodeGen/SPIRV/struct.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
 
 %struct.ST = type { i32, i32, i32 }
 
-; CHECK-SPIRV:     OpName %[[#struct:]] "struct.ST"
-; CHECK-SPIRV:     %[[#int:]] = OpTypeInt 32 0
-; CHECK-SPIRV-DAG: %[[#struct]] = OpTypeStruct %[[#int]] %[[#int]] %[[#int]]
-; CHECK-SPIRV-DAG: %[[#structP:]] = OpTypePointer Function %[[#struct]]
-; CHECK-SPIRV-DAG: %[[#intP:]] = OpTypePointer Function %[[#int]]
-; CHECK-SPIRV:     %[[#zero:]] = OpConstant %[[#int]] 0
-; CHECK-SPIRV:     %[[#one:]] = OpConstant %[[#int]] 1
-; CHECK-SPIRV:     %[[#two:]] = OpConstant %[[#int]] 2
-; CHECK-SPIRV:     %[[#three:]] = OpConstant %[[#int]] 3
+; CHECK-DAG: OpName %[[#struct:]] "struct.ST"
+; CHECK-DAG: %[[#char:]] = OpTypeInt 8 0
+; CHECK-DAG: %[[#int:]] = OpTypeInt 32 0
+; CHECK-DAG: %[[#struct]] = OpTypeStruct %[[#int]] %[[#int]] %[[#int]]
+; CHECK-DAG: %[[#structP:]] = OpTypePointer Function %[[#struct]]
+; CHECK-DAG: %[[#intP:]] = OpTypePointer Function %[[#char]]
+; CHECK-DAG: %[[#zero:]] = OpConstant %[[#int]] 0
+; CHECK-DAG: %[[#one:]] = OpConstant %[[#int]] 1
+; CHECK-DAG: %[[#two:]] = OpConstant %[[#int]] 2
+; CHECK-DAG: %[[#three:]] = OpConstant %[[#int]] 3
 
 define dso_local spir_func i32 @func() {
 entry:
 ; CHECK-SPIRV: %[[#st:]] = OpVariable %[[#structP]]
   %st = alloca %struct.ST, align 4
 ; CHECK-SPIRV: %[[#a:]] = OpInBoundsPtrAccessChain %[[#intP]] %[[#st]] %[[#zero]] %[[#zero]]
 ; CHECK-SPIRV: OpStore %[[#a]] %[[#one]]
   %a = getelementptr inbounds %struct.ST, %struct.ST* %st, i32 0, i32 0

llvm/test/CodeGen/SPIRV/transcoding/BuildNDRange_2.ll

 ;;   const size_t lsize2[2] = {1, 1};
 ;;   ndrange_2D(lsize2);
 ;;
 ;;   const size_t lsize3[3] = {1, 1, 1};
 ;;   ndrange_3D(lsize3);
 ;; }
 ;; bash$ $PATH_TO_GEN/bin/clang -cc1 -x cl -cl-std=CL2.0 -triple spir64-unknown-unknown -emit-llvm  -include opencl-20.h  BuildNDRange_2.cl -o BuildNDRange_2.ll
 
-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
 
 ; CHECK-SPIRV-DAG:     %[[#LEN2_ID:]] = OpConstant %[[#]] 2
 ; CHECK-SPIRV-DAG:     %[[#LEN3_ID:]] = OpConstant %[[#]] 3
 ; CHECK-SPIRV-DAG:     %[[#ARRAY_T2:]] = OpTypeArray %[[#]] %[[#LEN2_ID]]
 ; CHECK-SPIRV-DAG:     %[[#ARRAY_T3:]] = OpTypeArray %[[#]] %[[#LEN3_ID]]
 
 ; CHECK-SPIRV-LABEL:   OpFunction
 ; CHECK-SPIRV:         %[[#LOAD2_ID:]] = OpLoad %[[#ARRAY_T2]]

llvm/test/CodeGen/SPIRV/transcoding/OpConstantSampler.ll

 ;; 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 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; 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 {
+define spir_func <4 x float> @foo(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %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)
+  %0 = tail call target("spirv.Sampler") @__translate_sampler_initializer(i32 23)
+  %1 = tail call target("spirv.Sampler") @__translate_sampler_initializer(i32 0)
+  %call = tail call spir_func <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv2_ff(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %src, target("spirv.Sampler") %0, <2 x float> zeroinitializer, float 0.000000e+00)
+  %call1 = tail call spir_func <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv2_ff(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %src, target("spirv.Sampler") %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 target("spirv.Sampler") @__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
+declare spir_func <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv2_ff(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0), target("spirv.Sampler"), <2 x float>, float) local_unnamed_addr

llvm/test/CodeGen/SPIRV/transcoding/OpenCL/atomic_cmpxchg.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
 
 ;; This test checks that the backend is capable to correctly translate
 ;; atomic_cmpxchg OpenCL C 1.2 built-in function [1] into corresponding SPIR-V
 ;; instruction.
 
 ;; __kernel void test_atomic_cmpxchg(__global int *p, int cmp, int val) {
 ;;   atomic_cmpxchg(p, cmp, val);
 ;;
 ;;   __global unsigned int *up = (__global unsigned int *)p;
 ;;   unsigned int ucmp = (unsigned int)cmp;
 ;;   unsigned int uval = (unsigned int)val;
 ;;   atomic_cmpxchg(up, ucmp, uval);
 ;; }
 
 ; CHECK-SPIRV:     OpName %[[#TEST:]] "test_atomic_cmpxchg"
+; CHECK-SPIRV-DAG: %[[#UCHAR:]] = OpTypeInt 8 0
 ; CHECK-SPIRV-DAG: %[[#UINT:]] = OpTypeInt 32 0
-; CHECK-SPIRV-DAG: %[[#UINT_PTR:]] = OpTypePointer CrossWorkgroup %[[#UINT]]
+; CHECK-SPIRV-DAG: %[[#UCHAR_PTR:]] = OpTypePointer CrossWorkgroup %[[#UCHAR]]
 
 ;; In SPIR-V, atomic_cmpxchg is represented as OpAtomicCompareExchange [2],
 ;; which also includes memory scope and two memory semantic arguments. The
 ;; backend applies some default memory order for it and therefore, constants
 ;; below include a bit more information than original source
 
 ;; 0x2 Workgroup
 ; CHECK-SPIRV-DAG: %[[#WORKGROUP_SCOPE:]] = OpConstant %[[#UINT]] 2
 
 ;; 0x0 Relaxed
 ;; TODO: do we need CrossWorkgroupMemory here as well?
 ; CHECK-SPIRV-DAG: %[[#RELAXED:]] = OpConstant %[[#UINT]] 0
 
 ; CHECK-SPIRV:     %[[#TEST]] = OpFunction %[[#]]
-; CHECK-SPIRV:     %[[#PTR:]] = OpFunctionParameter %[[#UINT_PTR]]
+; CHECK-SPIRV:     %[[#PTR:]] = OpFunctionParameter %[[#UCHAR_PTR]]
 ; CHECK-SPIRV:     %[[#CMP:]] = OpFunctionParameter %[[#UINT]]
 ; CHECK-SPIRV:     %[[#VAL:]] = OpFunctionParameter %[[#UINT]]
 ; CHECK-SPIRV:     %[[#]] = OpAtomicCompareExchange %[[#UINT]] %[[#PTR]] %[[#WORKGROUP_SCOPE]] %[[#RELAXED]] %[[#RELAXED]] %[[#VAL]] %[[#CMP]]
 ; CHECK-SPIRV:     %[[#]] = OpAtomicCompareExchange %[[#UINT]] %[[#PTR]] %[[#WORKGROUP_SCOPE]] %[[#RELAXED]] %[[#RELAXED]] %[[#VAL]] %[[#CMP]]
 
 define dso_local spir_kernel void @test_atomic_cmpxchg(i32 addrspace(1)* noundef %p, i32 noundef %cmp, i32 noundef %val) local_unnamed_addr {
 entry:
   %call = tail call spir_func i32 @_Z14atomic_cmpxchgPU3AS1Viii(i32 addrspace(1)* noundef %p, i32 noundef %cmp, i32 noundef %val)

llvm/test/CodeGen/SPIRV/transcoding/OpenCL/atomic_legacy.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
 
 ;; This test checks that the backend is capable to correctly translate
 ;; legacy atomic OpenCL C 1.2 built-in functions [1] into corresponding SPIR-V
 ;; instruction.
 
 ;; __kernel void test_legacy_atomics(__global int *p, int val) {
 ;;   atom_add(p, val);     // from cl_khr_global_int32_base_atomics
 ;;   atomic_add(p, val);   // from OpenCL C 1.1
 ;; }
 
 ; CHECK-SPIRV:     OpName %[[#TEST:]] "test_legacy_atomics"
 ; CHECK-SPIRV-DAG: %[[#UINT:]] = OpTypeInt 32 0
-; CHECK-SPIRV-DAG: %[[#UINT_PTR:]] = OpTypePointer CrossWorkgroup %[[#UINT]]
+; CHECK-SPIRV-DAG: %[[#UCHAR:]] = OpTypeInt 8 0
+; CHECK-SPIRV-DAG: %[[#UCHAR_PTR:]] = OpTypePointer CrossWorkgroup %[[#UCHAR]]
 
 ;; In SPIR-V, atomic_add is represented as OpAtomicIAdd [2], which also includes
 ;; memory scope and memory semantic arguments. The backend applies a default
 ;; memory scope and memory order for it and therefore, constants below include
 ;; a bit more information than original source
 
 ;; 0x2 Workgroup
 ; CHECK-SPIRV-DAG: %[[#WORKGROUP_SCOPE:]] = OpConstant %[[#UINT]] 2
 
 ;; 0x0 Relaxed
 ; CHECK-SPIRV-DAG: %[[#RELAXED:]] = OpConstant %[[#UINT]] 0
 
 ; CHECK-SPIRV:     %[[#TEST]] = OpFunction %[[#]]
-; CHECK-SPIRV:     %[[#PTR:]] = OpFunctionParameter %[[#UINT_PTR]]
+; CHECK-SPIRV:     %[[#PTR:]] = OpFunctionParameter %[[#UCHAR_PTR]]
 ; CHECK-SPIRV:     %[[#VAL:]] = OpFunctionParameter %[[#UINT]]
 ; CHECK-SPIRV:     %[[#]] = OpAtomicIAdd %[[#UINT]] %[[#PTR]] %[[#WORKGROUP_SCOPE]] %[[#RELAXED]] %[[#VAL]]
 ; CHECK-SPIRV:     %[[#]] = OpAtomicIAdd %[[#UINT]] %[[#PTR]] %[[#WORKGROUP_SCOPE]] %[[#RELAXED]] %[[#VAL]]
 
 define dso_local spir_kernel void @test_legacy_atomics(i32 addrspace(1)* noundef %p, i32 noundef %val) local_unnamed_addr {
 entry:
   %call = tail call spir_func i32 @_Z8atom_addPU3AS1Vii(i32 addrspace(1)* noundef %p, i32 noundef %val)
   %call1 = tail call spir_func i32 @_Z10atomic_addPU3AS1Vii(i32 addrspace(1)* noundef %p, i32 noundef %val)

llvm/test/CodeGen/SPIRV/transcoding/block_w_struct_return.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefixes=CHECK-SPIRV,CHECK-SPIRV1_4
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefixes=CHECK-SPIRV,CHECK-SPIRV1_4
 
 ;; TODO: We cannot check SPIR_V 1.1 and 1.4 simultaneously, implement additional
 ;;       run with CHECK-SPIRV1_1.
 
 ;; kernel void block_ret_struct(__global int* res)
 ;; {
 ;;   struct A {
 ;;       int a;

llvm/test/CodeGen/SPIRV/transcoding/enqueue_kernel.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
 
 ; CHECK-SPIRV: OpEntryPoint Kernel %[[#BlockKer1:]] "__device_side_enqueue_block_invoke_kernel"
 ; CHECK-SPIRV: OpEntryPoint Kernel %[[#BlockKer2:]] "__device_side_enqueue_block_invoke_2_kernel"
 ; CHECK-SPIRV: OpEntryPoint Kernel %[[#BlockKer3:]] "__device_side_enqueue_block_invoke_3_kernel"
 ; CHECK-SPIRV: OpEntryPoint Kernel %[[#BlockKer4:]] "__device_side_enqueue_block_invoke_4_kernel"
 ; CHECK-SPIRV: OpEntryPoint Kernel %[[#BlockKer5:]] "__device_side_enqueue_block_invoke_5_kernel"
 ; CHECK-SPIRV: OpName %[[#BlockGlb1:]] "__block_literal_global"
 ; CHECK-SPIRV: OpName %[[#BlockGlb2:]] "__block_literal_global.1"

llvm/test/CodeGen/SPIRV/transcoding/extract_insert_value.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
-
-;; Check 'LLVM ==> SPIR-V' conversion of extractvalue/insertvalue.
-
-%struct.arr = type { [7 x float] }
-%struct.st = type { %struct.inner }
-%struct.inner = type { float }
-
-; CHECK-SPIRV:     %[[#float_ty:]] = OpTypeFloat 32
-; CHECK-SPIRV:     %[[#int_ty:]] = OpTypeInt 32
-; CHECK-SPIRV:     %[[#arr_size:]] = OpConstant %[[#int_ty]] 7
-; CHECK-SPIRV:     %[[#array_ty:]] = OpTypeArray %[[#float_ty]] %[[#arr_size]]
-; CHECK-SPIRV:     %[[#struct_ty:]] = OpTypeStruct %[[#array_ty]]
-; CHECK-SPIRV:     %[[#struct_ptr_ty:]] = OpTypePointer CrossWorkgroup %[[#struct_ty]]
-; CHECK-SPIRV:     %[[#array_ptr_ty:]] = OpTypePointer CrossWorkgroup %[[#array_ty]]
-; CHECK-SPIRV:     %[[#struct1_in_ty:]] = OpTypeStruct %[[#float_ty]]
-; CHECK-SPIRV:     %[[#struct1_ty:]] = OpTypeStruct %[[#struct1_in_ty]]
-; CHECK-SPIRV:     %[[#struct1_ptr_ty:]] = OpTypePointer CrossWorkgroup %[[#struct1_ty]]
-; CHECK-SPIRV:     %[[#struct1_in_ptr_ty:]] = OpTypePointer CrossWorkgroup %[[#struct1_in_ty]]
-; CHECK-SPIRV-NOT: OpConstant %{{.*}} 2
-; CHECK-SPIRV-NOT: OpConstant %{{.*}} 4
-; CHECK-SPIRV-NOT: OpConstant %{{.*}} 5
-
-; CHECK-SPIRV-LABEL:  OpFunction
-; CHECK-SPIRV-NEXT:   %[[#object:]] = OpFunctionParameter %[[#struct_ptr_ty]]
-; CHECK-SPIRV:        %[[#store_ptr:]] = OpInBoundsPtrAccessChain %[[#array_ptr_ty]] %[[#object]] %[[#]] %[[#]]
-; CHECK-SPIRV:        %[[#extracted_array:]] = OpLoad %[[#array_ty]] %[[#store_ptr]] Aligned 4
-; CHECK-SPIRV:        %[[#elem_4:]] = OpCompositeExtract %[[#float_ty]] %[[#extracted_array]] 4
-; CHECK-SPIRV:        %[[#elem_2:]] = OpCompositeExtract %[[#float_ty]] %[[#extracted_array]] 2
-; CHECK-SPIRV:        %[[#add:]] = OpFAdd %[[#float_ty]] %[[#elem_4]] %[[#elem_2]]
-; CHECK-SPIRV:        %[[#inserted_array:]] = OpCompositeInsert %[[#array_ty]] %[[#add]] %[[#extracted_array]] 5
-; CHECK-SPIRV:        OpStore %[[#store_ptr]] %[[#inserted_array]]
-; CHECK-SPIRV-LABEL:  OpFunctionEnd
-
-define spir_func void @array_test(%struct.arr addrspace(1)* %object) {
-entry:
-  %0 = getelementptr inbounds %struct.arr, %struct.arr addrspace(1)* %object, i32 0, i32 0
-  %1 = load [7 x float], [7 x float] addrspace(1)* %0, align 4
-  %2 = extractvalue [7 x float] %1, 4
-  %3 = extractvalue [7 x float] %1, 2
-  %4 = fadd float %2, %3
-  %5 = insertvalue [7 x float] %1, float %4, 5
-  store [7 x float] %5, [7 x float] addrspace(1)* %0
-  ret void
-}
-
-; CHECK-SPIRV-LABEL:  OpFunction
-; CHECK-SPIRV-NEXT:   %[[#object:]] = OpFunctionParameter %[[#struct1_ptr_ty]]
-; CHECK-SPIRV:        %[[#store1_ptr:]] = OpInBoundsPtrAccessChain %[[#struct1_in_ptr_ty]] %[[#object]] %[[#]] %[[#]]
-; CHECK-SPIRV:        %[[#extracted_struct:]] = OpLoad %[[#struct1_in_ty]] %[[#store1_ptr]] Aligned 4
-; CHECK-SPIRV:        %[[#elem:]] = OpCompositeExtract %[[#float_ty]] %[[#extracted_struct]] 0
-; CHECK-SPIRV:        %[[#add:]] = OpFAdd %[[#float_ty]] %[[#elem]] %[[#]]
-; CHECK-SPIRV:        %[[#inserted_struct:]] = OpCompositeInsert %[[#struct1_in_ty]] %[[#add]] %[[#extracted_struct]] 0
-; CHECK-SPIRV:        OpStore %[[#store1_ptr]] %[[#inserted_struct]]
-; CHECK-SPIRV-LABEL:  OpFunctionEnd
-
-define spir_func void @struct_test(%struct.st addrspace(1)* %object) {
-entry:
-  %0 = getelementptr inbounds %struct.st, %struct.st addrspace(1)* %object, i32 0, i32 0
-  %1 = load %struct.inner, %struct.inner addrspace(1)* %0, align 4
-  %2 = extractvalue %struct.inner %1, 0
-  %3 = fadd float %2, 1.000000e+00
-  %4 = insertvalue %struct.inner %1, float %3, 0
-  store %struct.inner %4, %struct.inner addrspace(1)* %0
-  ret void
-}

llvm/test/CodeGen/SPIRV/transcoding/get_image_num_mip_levels.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
-
-;; Types:
-; CHECK-DAG:  %[[#INT:]] = OpTypeInt 32
-; CHECK-DAG:  %[[#VOID:]] = OpTypeVoid
-; CHECK-DAG:  %[[#IMAGE1D_T:]] = OpTypeImage %[[#VOID]] 1D 0 0 0 0 Unknown ReadOnly
-; CHECK-DAG:  %[[#IMAGE2D_T:]] = OpTypeImage %[[#VOID]] 2D 0 0 0 0 Unknown ReadOnly
-; CHECK-DAG:  %[[#IMAGE3D_T:]] = OpTypeImage %[[#VOID]] 3D 0 0 0 0 Unknown ReadOnly
-; CHECK-DAG:  %[[#IMAGE1D_ARRAY_T:]] = OpTypeImage %[[#VOID]] 1D 0 1 0 0 Unknown ReadOnly
-; CHECK-DAG:  %[[#IMAGE2D_ARRAY_T:]] = OpTypeImage %[[#VOID]] 2D 0 1 0 0 Unknown ReadOnly
-; CHECK-DAG:  %[[#IMAGE2D_DEPTH_T:]] = OpTypeImage %[[#VOID]] 2D 1 0 0 0 Unknown ReadOnly
-; CHECK-DAG:  %[[#IMAGE2D_ARRAY_DEPTH_T:]] = OpTypeImage %[[#VOID]] 2D 1 1 0 0 Unknown ReadOnly
-;; Instructions:
-; CHECK:      %[[#IMAGE1D:]] = OpLoad %[[#IMAGE1D_T]]
-; CHECK-NEXT: %[[#]] = OpImageQueryLevels %[[#INT]] %[[#IMAGE1D]]
-; CHECK:      %[[#IMAGE2D:]] = OpLoad %[[#IMAGE2D_T]]
-; CHECK-NEXT: %[[#]] = OpImageQueryLevels %[[#INT]] %[[#IMAGE2D]]
-; CHECK:      %[[#IMAGE3D:]] = OpLoad %[[#IMAGE3D_T]]
-; CHECK-NEXT: %[[#]] = OpImageQueryLevels %[[#INT]] %[[#IMAGE3D]]
-; CHECK:      %[[#IMAGE1D_ARRAY:]] = OpLoad %[[#IMAGE1D_ARRAY_T]]
-; CHECK-NEXT: %[[#]] = OpImageQueryLevels %[[#INT]] %[[#IMAGE1D_ARRAY]]
-; CHECK:      %[[#IMAGE2D_ARRAY:]] = OpLoad %[[#IMAGE2D_ARRAY_T]]
-; CHECK-NEXT: %[[#]] = OpImageQueryLevels %[[#INT]] %[[#IMAGE2D_ARRAY]]
-; CHECK:      %[[#IMAGE2D_DEPTH:]] = OpLoad %[[#IMAGE2D_DEPTH_T]]
-; CHECK-NEXT: %[[#]] = OpImageQueryLevels %[[#INT]] %[[#IMAGE2D_DEPTH]]
-; CHECK:      %[[#IMAGE2D_ARRAY_DEPTH:]] = OpLoad %[[#IMAGE2D_ARRAY_DEPTH_T]]
-; CHECK-NEXT: %[[#]] = OpImageQueryLevels %[[#INT]] %[[#IMAGE2D_ARRAY_DEPTH]]
-
-%opencl.image1d_ro_t = type opaque
-%opencl.image2d_ro_t = type opaque
-%opencl.image3d_ro_t = type opaque
-%opencl.image1d_array_ro_t = type opaque
-%opencl.image2d_array_ro_t = type opaque
-%opencl.image2d_depth_ro_t = type opaque
-%opencl.image2d_array_depth_ro_t = type opaque
-
-define spir_func void @testimage1d(%opencl.image1d_ro_t addrspace(1)* %img1, %opencl.image2d_ro_t addrspace(1)* %img2, %opencl.image3d_ro_t addrspace(1)* %img3, %opencl.image1d_array_ro_t addrspace(1)* %img4, %opencl.image2d_array_ro_t addrspace(1)* %img5, %opencl.image2d_depth_ro_t addrspace(1)* %img6, %opencl.image2d_array_depth_ro_t addrspace(1)* %img7) {
-entry:
-  %img1.addr = alloca %opencl.image1d_ro_t addrspace(1)*, align 4
-  %img2.addr = alloca %opencl.image2d_ro_t addrspace(1)*, align 4
-  %img3.addr = alloca %opencl.image3d_ro_t addrspace(1)*, align 4
-  %img4.addr = alloca %opencl.image1d_array_ro_t addrspace(1)*, align 4
-  %img5.addr = alloca %opencl.image2d_array_ro_t addrspace(1)*, align 4
-  %img6.addr = alloca %opencl.image2d_depth_ro_t addrspace(1)*, align 4
-  %img7.addr = alloca %opencl.image2d_array_depth_ro_t addrspace(1)*, align 4
-  store %opencl.image1d_ro_t addrspace(1)* %img1, %opencl.image1d_ro_t addrspace(1)** %img1.addr, align 4
-  store %opencl.image2d_ro_t addrspace(1)* %img2, %opencl.image2d_ro_t addrspace(1)** %img2.addr, align 4
-  store %opencl.image3d_ro_t addrspace(1)* %img3, %opencl.image3d_ro_t addrspace(1)** %img3.addr, align 4
-  store %opencl.image1d_array_ro_t addrspace(1)* %img4, %opencl.image1d_array_ro_t addrspace(1)** %img4.addr, align 4
-  store %opencl.image2d_array_ro_t addrspace(1)* %img5, %opencl.image2d_array_ro_t addrspace(1)** %img5.addr, align 4
-  store %opencl.image2d_depth_ro_t addrspace(1)* %img6, %opencl.image2d_depth_ro_t addrspace(1)** %img6.addr, align 4
-  store %opencl.image2d_array_depth_ro_t addrspace(1)* %img7, %opencl.image2d_array_depth_ro_t addrspace(1)** %img7.addr, align 4
-  %0 = load %opencl.image1d_ro_t addrspace(1)*, %opencl.image1d_ro_t addrspace(1)** %img1.addr, align 4
-  %call = call spir_func i32 @_Z24get_image_num_mip_levels14ocl_image1d_ro(%opencl.image1d_ro_t addrspace(1)* %0)
-  %1 = load %opencl.image2d_ro_t addrspace(1)*, %opencl.image2d_ro_t addrspace(1)** %img2.addr, align 4
-  %call1 = call spir_func i32 @_Z24get_image_num_mip_levels14ocl_image2d_ro(%opencl.image2d_ro_t addrspace(1)* %1)
-  %2 = load %opencl.image3d_ro_t addrspace(1)*, %opencl.image3d_ro_t addrspace(1)** %img3.addr, align 4
-  %call2 = call spir_func i32 @_Z24get_image_num_mip_levels14ocl_image3d_ro(%opencl.image3d_ro_t addrspace(1)* %2)
-  %3 = load %opencl.image1d_array_ro_t addrspace(1)*, %opencl.image1d_array_ro_t addrspace(1)** %img4.addr, align 4
-  %call3 = call spir_func i32 @_Z24get_image_num_mip_levels20ocl_image1d_array_ro(%opencl.image1d_array_ro_t addrspace(1)* %3)
-  %4 = load %opencl.image2d_array_ro_t addrspace(1)*, %opencl.image2d_array_ro_t addrspace(1)** %img5.addr, align 4
-  %call4 = call spir_func i32 @_Z24get_image_num_mip_levels20ocl_image2d_array_ro(%opencl.image2d_array_ro_t addrspace(1)* %4)
-  %5 = load %opencl.image2d_depth_ro_t addrspace(1)*, %opencl.image2d_depth_ro_t addrspace(1)** %img6.addr, align 4
-  %call5 = call spir_func i32 @_Z24get_image_num_mip_levels20ocl_image2d_depth_ro(%opencl.image2d_depth_ro_t addrspace(1)* %5)
-  %6 = load %opencl.image2d_array_depth_ro_t addrspace(1)*, %opencl.image2d_array_depth_ro_t addrspace(1)** %img7.addr, align 4
-  %call6 = call spir_func i32 @_Z24get_image_num_mip_levels26ocl_image2d_array_depth_ro(%opencl.image2d_array_depth_ro_t addrspace(1)* %6)
-  ret void
-}
-
-declare spir_func i32 @_Z24get_image_num_mip_levels14ocl_image1d_ro(%opencl.image1d_ro_t addrspace(1)*)
-
-declare spir_func i32 @_Z24get_image_num_mip_levels14ocl_image2d_ro(%opencl.image2d_ro_t addrspace(1)*)
-
-declare spir_func i32 @_Z24get_image_num_mip_levels14ocl_image3d_ro(%opencl.image3d_ro_t addrspace(1)*)
-
-declare spir_func i32 @_Z24get_image_num_mip_levels20ocl_image1d_array_ro(%opencl.image1d_array_ro_t addrspace(1)*)
-
-declare spir_func i32 @_Z24get_image_num_mip_levels20ocl_image2d_array_ro(%opencl.image2d_array_ro_t addrspace(1)*)
-
-declare spir_func i32 @_Z24get_image_num_mip_levels20ocl_image2d_depth_ro(%opencl.image2d_depth_ro_t addrspace(1)*)
-
-declare spir_func i32 @_Z24get_image_num_mip_levels26ocl_image2d_array_depth_ro(%opencl.image2d_array_depth_ro_t addrspace(1)*)

llvm/test/CodeGen/SPIRV/transcoding/global_block.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefixes=CHECK-SPIRV,CHECK-SPIRV1_4
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefixes=CHECK-SPIRV,CHECK-SPIRV1_4
 
 ;; There are no blocks in SPIR-V. Therefore they are translated into regular
 ;; functions. An LLVM module which uses blocks, also contains some auxiliary
 ;; block-specific instructions, which are redundant in SPIR-V and should be
 ;; removed
 
 ;; kernel void block_kernel(__global int* res) {
 ;;   typedef int (^block_t)(int);

llvm/test/CodeGen/SPIRV/transcoding/image_with_access_qualifiers.ll

-; RUN: llc -O0 -opaque-pointers=0 -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)

llvm/test/CodeGen/SPIRV/transcoding/spec_const.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
 
 ; CHECK-SPIRV-NOT: OpCapability Matrix
 ; CHECK-SPIRV-NOT: OpCapability Shader
-; CHECK-SPIRV:     OpCapability Float16Buffer
+; CHECK-SPIRV:     OpCapability Kernel
 
 ; CHECK-SPIRV-DAG: OpDecorate %[[#SC0:]] SpecId 0
 ; CHECK-SPIRV-DAG: OpDecorate %[[#SC1:]] SpecId 1
 ; CHECK-SPIRV-DAG: OpDecorate %[[#SC2:]] SpecId 2
 ; CHECK-SPIRV-DAG: OpDecorate %[[#SC3:]] SpecId 3
 ; CHECK-SPIRV-DAG: OpDecorate %[[#SC4:]] SpecId 4
 ; CHECK-SPIRV-DAG: OpDecorate %[[#SC5:]] SpecId 5
 ; CHECK-SPIRV-DAG: OpDecorate %[[#SC6:]] SpecId 6

llvm/test/CodeGen/SPIRV/transcoding/spirv-private-array-initialization.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
 ;
-; CHECK-SPIRV-DAG: %[[#i32:]] = OpTypeInt 32 0
 ; CHECK-SPIRV-DAG: %[[#i8:]] = OpTypeInt 8 0
+; CHECK-SPIRV-DAG: %[[#i32:]] = OpTypeInt 32 0
 ; CHECK-SPIRV-DAG: %[[#one:]] = OpConstant %[[#i32]] 1
 ; CHECK-SPIRV-DAG: %[[#two:]] = OpConstant %[[#i32]] 2
 ; CHECK-SPIRV-DAG: %[[#three:]] = OpConstant %[[#i32]] 3
 ; CHECK-SPIRV-DAG: %[[#i32x3:]] = OpTypeArray %[[#i32]] %[[#three]]
-; CHECK-SPIRV-DAG: %[[#i32x3_ptr:]] = OpTypePointer Function %[[#i32x3]]
+; CHECK-SPIRV-DAG: %[[#test_arr_init:]] = OpConstantComposite %[[#i32x3]] %[[#one]] %[[#two]] %[[#three]]
+; CHECK-SPIRV-DAG: %[[#twelve:]] = OpConstant %[[#i32]] 12
 ; CHECK-SPIRV-DAG: %[[#const_i32x3_ptr:]] = OpTypePointer UniformConstant %[[#i32x3]]
-; CHECK-SPIRV-DAG: %[[#i8_ptr:]] = OpTypePointer Function %[[#i8]]
-; CHECK-SPIRV-DAG: %[[#const_i8_ptr:]] = OpTypePointer UniformConstant %[[#i8]]
-; CHECK-SPIRV:     %[[#test_arr_init:]] = OpConstantComposite %[[#i32x3]] %[[#one]] %[[#two]] %[[#three]]
-; CHECK-SPIRV:     %[[#twelve:]] = OpConstant %[[#i32]] 12
+
 ; CHECK-SPIRV:     %[[#test_arr2:]] = OpVariable %[[#const_i32x3_ptr]] UniformConstant %[[#test_arr_init]]
 ; CHECK-SPIRV:     %[[#test_arr:]] = OpVariable %[[#const_i32x3_ptr]] UniformConstant %[[#test_arr_init]]
-;
+
+; CHECK-SPIRV-DAG: %[[#i8_ptr:]] = OpTypePointer Function %[[#i8]]
+; CHECK-SPIRV-DAG: %[[#const_i8_ptr:]] = OpTypePointer UniformConstant %[[#i8]]
+; CHECK-SPIRV-DAG: %[[#i32x3_ptr:]] = OpTypePointer Function %[[#i32x3]]
+
 ; CHECK-SPIRV:     %[[#arr:]] = OpVariable %[[#i32x3_ptr]] Function
 ; CHECK-SPIRV:     %[[#arr2:]] = OpVariable %[[#i32x3_ptr]] Function
-;
 ; CHECK-SPIRV:     %[[#arr_i8_ptr:]] = OpBitcast %[[#i8_ptr]] %[[#arr]]
-; CHECK-SPIRV:     %[[#test_arr_const_i8_ptr:]] = OpBitcast %[[#const_i8_ptr]] %[[#test_arr]]
-; CHECK-SPIRV:     OpCopyMemorySized %[[#arr_i8_ptr]] %[[#test_arr_const_i8_ptr]] %[[#twelve]] Aligned 4
-;
+; CHECK-SPIRV:     OpCopyMemorySized %[[#arr_i8_ptr]] %[[#test_arr]] %[[#twelve]] Aligned 4
 ; CHECK-SPIRV:     %[[#arr2_i8_ptr:]] = OpBitcast %[[#i8_ptr]] %[[#arr2]]
-; CHECK-SPIRV:     %[[#test_arr2_const_i8_ptr:]] = OpBitcast %[[#const_i8_ptr]] %[[#test_arr2]]
-; CHECK-SPIRV:     OpCopyMemorySized %[[#arr2_i8_ptr]] %[[#test_arr2_const_i8_ptr]] %[[#twelve]] Aligned 4
+; CHECK-SPIRV:     OpCopyMemorySized %[[#arr2_i8_ptr]] %[[#test_arr2]] %[[#twelve]] Aligned 4
+
 
 @__const.test.arr = private unnamed_addr addrspace(2) constant [3 x i32] [i32 1, i32 2, i32 3], align 4
 @__const.test.arr2 = private unnamed_addr addrspace(2) constant [3 x i32] [i32 1, i32 2, i32 3], align 4
 
 define spir_func void @test() {
 entry:
   %arr = alloca [3 x i32], align 4
   %arr2 = alloca [3 x i32], align 4
   %0 = bitcast [3 x i32]* %arr to i8*
   call void @llvm.memcpy.p0i8.p2i8.i32(i8* align 4 %0, i8 addrspace(2)* align 4 bitcast ([3 x i32] addrspace(2)* @__const.test.arr to i8 addrspace(2)*), i32 12, i1 false)
   %1 = bitcast [3 x i32]* %arr2 to i8*
   call void @llvm.memcpy.p0i8.p2i8.i32(i8* align 4 %1, i8 addrspace(2)* align 4 bitcast ([3 x i32] addrspace(2)* @__const.test.arr2 to i8 addrspace(2)*), i32 12, i1 false)
   ret void
 }
 
 declare void @llvm.memcpy.p0i8.p2i8.i32(i8* nocapture writeonly, i8 addrspace(2)* nocapture readonly, i32, i1)

llvm/test/CodeGen/SPIRV/types/device_event.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; 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)
+  %call = call spir_func target("spirv.DeviceEvent") @_Z17create_user_eventv()
+  %call1 = call spir_func zeroext i1 @_Z14is_valid_event12ocl_clkevent(target("spirv.DeviceEvent") %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)
+  call spir_func void @_Z12retain_event12ocl_clkevent(target("spirv.DeviceEvent") %call)
+  call spir_func void @_Z21set_user_event_status12ocl_clkeventi(target("spirv.DeviceEvent") %call, i32 noundef -42)
+  call spir_func void @_Z28capture_event_profiling_info12ocl_clkeventiPU3AS1v(target("spirv.DeviceEvent") %call, i32 noundef 1, i8 addrspace(1)* noundef %prof)
+  call spir_func void @_Z13release_event12ocl_clkevent(target("spirv.DeviceEvent") %call)
   ret void
 }
 
-declare spir_func %opencl.clk_event_t* @_Z17create_user_eventv() local_unnamed_addr
+declare spir_func target("spirv.DeviceEvent") @_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 zeroext i1 @_Z14is_valid_event12ocl_clkevent(target("spirv.DeviceEvent")) local_unnamed_addr
 
-declare spir_func void @_Z12retain_event12ocl_clkevent(%opencl.clk_event_t*) local_unnamed_addr
+declare spir_func void @_Z12retain_event12ocl_clkevent(target("spirv.DeviceEvent")) 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 @_Z21set_user_event_status12ocl_clkeventi(target("spirv.DeviceEvent"), 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 @_Z28capture_event_profiling_info12ocl_clkeventiPU3AS1v(target("spirv.DeviceEvent"), i32 noundef, i8 addrspace(1)* noundef) local_unnamed_addr
 
-declare spir_func void @_Z13release_event12ocl_clkevent(%opencl.clk_event_t*) local_unnamed_addr
+declare spir_func void @_Z13release_event12ocl_clkevent(target("spirv.DeviceEvent")) local_unnamed_addr

llvm/test/CodeGen/SPIRV/types/image-access-qualifier.ll

+; 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]]
+
+define spir_func void @sampFun(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 2) %image) {
+entry:
+  %image.addr = alloca target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 2), align 4
+  store target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 2) %image, ptr %image.addr, align 4
+  %0 = load target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 2), ptr %image.addr, align 4
+  %call = call spir_func <4 x float> @_Z11read_imagef14ocl_image1d_rw11ocl_sampleri(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 2) %0, i32 8, i32 2)
+  ret void
+}
+
+declare spir_func <4 x float> @_Z11read_imagef14ocl_image1d_rw11ocl_sampleri(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 2), i32, i32)

llvm/test/CodeGen/SPIRV/types/image-function-argument.ll

+; 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]]
+
+define spir_func void @f0(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %v2, <2 x float> %v3) {
+  ret void
+}
+
+define spir_func void @f1(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %v2, <2 x float> %v3) {
+  ret void
+}
+
+define spir_kernel void @test(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %v1) {
+  call spir_func void @f0(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %v1, <2 x float> <float 1.000000e+00, float 5.000000e+00>)
+  ret void
+}

llvm/test/CodeGen/SPIRV/types/image-query-dim-array-size.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; 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) {
+define spir_kernel void @sample_kernel(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0) %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.tmp1 = call spir_func <2 x i32> @_Z13get_image_dim20ocl_image2d_array_ro(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0) %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.tmp3 = call spir_func i64 @_Z20get_image_array_size20ocl_image2d_array_ro(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0) %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 <2 x i32> @_Z13get_image_dim20ocl_image2d_array_ro(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0))
 
-declare spir_func i64 @_Z20get_image_array_size20ocl_image2d_array_ro(%opencl.image2d_array_ro_t addrspace(1)*)
+declare spir_func i64 @_Z20get_image_array_size20ocl_image2d_array_ro(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0))

llvm/test/CodeGen/SPIRV/types/image-query-mip-levels.ll

+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+
+;; Types:
+; CHECK-DAG:  %[[#INT:]] = OpTypeInt 32
+; CHECK-DAG:  %[[#VOID:]] = OpTypeVoid
+; CHECK-DAG:  %[[#IMAGE1D_T:]] = OpTypeImage %[[#VOID]] 1D 0 0 0 0 Unknown ReadOnly
+; CHECK-DAG:  %[[#IMAGE2D_T:]] = OpTypeImage %[[#VOID]] 2D 0 0 0 0 Unknown ReadOnly
+; CHECK-DAG:  %[[#IMAGE3D_T:]] = OpTypeImage %[[#VOID]] 3D 0 0 0 0 Unknown ReadOnly
+; CHECK-DAG:  %[[#IMAGE1D_ARRAY_T:]] = OpTypeImage %[[#VOID]] 1D 0 1 0 0 Unknown ReadOnly
+; CHECK-DAG:  %[[#IMAGE2D_ARRAY_T:]] = OpTypeImage %[[#VOID]] 2D 0 1 0 0 Unknown ReadOnly
+; CHECK-DAG:  %[[#IMAGE2D_DEPTH_T:]] = OpTypeImage %[[#VOID]] 2D 1 0 0 0 Unknown ReadOnly
+; CHECK-DAG:  %[[#IMAGE2D_ARRAY_DEPTH_T:]] = OpTypeImage %[[#VOID]] 2D 1 1 0 0 Unknown ReadOnly
+;; Instructions:
+; CHECK:      %[[#IMAGE1D:]] = OpLoad %[[#IMAGE1D_T]]
+; CHECK-NEXT: %[[#]] = OpImageQueryLevels %[[#INT]] %[[#IMAGE1D]]
+; CHECK:      %[[#IMAGE2D:]] = OpLoad %[[#IMAGE2D_T]]
+; CHECK-NEXT: %[[#]] = OpImageQueryLevels %[[#INT]] %[[#IMAGE2D]]
+; CHECK:      %[[#IMAGE3D:]] = OpLoad %[[#IMAGE3D_T]]
+; CHECK-NEXT: %[[#]] = OpImageQueryLevels %[[#INT]] %[[#IMAGE3D]]
+; CHECK:      %[[#IMAGE1D_ARRAY:]] = OpLoad %[[#IMAGE1D_ARRAY_T]]
+; CHECK-NEXT: %[[#]] = OpImageQueryLevels %[[#INT]] %[[#IMAGE1D_ARRAY]]
+; CHECK:      %[[#IMAGE2D_ARRAY:]] = OpLoad %[[#IMAGE2D_ARRAY_T]]
+; CHECK-NEXT: %[[#]] = OpImageQueryLevels %[[#INT]] %[[#IMAGE2D_ARRAY]]
+; CHECK:      %[[#IMAGE2D_DEPTH:]] = OpLoad %[[#IMAGE2D_DEPTH_T]]
+; CHECK-NEXT: %[[#]] = OpImageQueryLevels %[[#INT]] %[[#IMAGE2D_DEPTH]]
+; CHECK:      %[[#IMAGE2D_ARRAY_DEPTH:]] = OpLoad %[[#IMAGE2D_ARRAY_DEPTH_T]]
+; CHECK-NEXT: %[[#]] = OpImageQueryLevels %[[#INT]] %[[#IMAGE2D_ARRAY_DEPTH]]
+
+define spir_func void @testimage1d(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0) %img1, target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %img2, target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %img3, target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 0) %img4, target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0) %img5, target("spirv.Image", void, 1, 1, 0, 0, 0, 0, 0) %img6, target("spirv.Image", void, 1, 1, 1, 0, 0, 0, 0) %img7) {
+entry:
+  %img1.addr = alloca target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0), align 4
+  %img2.addr = alloca target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0), align 4
+  %img3.addr = alloca target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0), align 4
+  %img4.addr = alloca target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 0), align 4
+  %img5.addr = alloca target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0), align 4
+  %img6.addr = alloca target("spirv.Image", void, 1, 1, 0, 0, 0, 0, 0), align 4
+  %img7.addr = alloca target("spirv.Image", void, 1, 1, 1, 0, 0, 0, 0), align 4
+  store target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0) %img1, target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0)* %img1.addr, align 4
+  store target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %img2, target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0)* %img2.addr, align 4
+  store target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %img3, target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0)* %img3.addr, align 4
+  store target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 0) %img4, target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 0)* %img4.addr, align 4
+  store target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0) %img5, target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0)* %img5.addr, align 4
+  store target("spirv.Image", void, 1, 1, 0, 0, 0, 0, 0) %img6, target("spirv.Image", void, 1, 1, 0, 0, 0, 0, 0)* %img6.addr, align 4
+  store target("spirv.Image", void, 1, 1, 1, 0, 0, 0, 0) %img7, target("spirv.Image", void, 1, 1, 1, 0, 0, 0, 0)* %img7.addr, align 4
+  %0 = load target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0), target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0)* %img1.addr, align 4
+  %call = call spir_func i32 @_Z24get_image_num_mip_levels14ocl_image1d_ro(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0) %0)
+  %1 = load target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0), target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0)* %img2.addr, align 4
+  %call1 = call spir_func i32 @_Z24get_image_num_mip_levels14ocl_image2d_ro(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %1)
+  %2 = load target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0), target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0)* %img3.addr, align 4
+  %call2 = call spir_func i32 @_Z24get_image_num_mip_levels14ocl_image3d_ro(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %2)
+  %3 = load target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 0), target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 0)* %img4.addr, align 4
+  %call3 = call spir_func i32 @_Z24get_image_num_mip_levels20ocl_image1d_array_ro(target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 0) %3)
+  %4 = load target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0), target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0)* %img5.addr, align 4
+  %call4 = call spir_func i32 @_Z24get_image_num_mip_levels20ocl_image2d_array_ro(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0) %4)
+  %5 = load target("spirv.Image", void, 1, 1, 0, 0, 0, 0, 0), target("spirv.Image", void, 1, 1, 0, 0, 0, 0, 0)* %img6.addr, align 4
+  %call5 = call spir_func i32 @_Z24get_image_num_mip_levels20ocl_image2d_depth_ro(target("spirv.Image", void, 1, 1, 0, 0, 0, 0, 0) %5)
+  %6 = load target("spirv.Image", void, 1, 1, 1, 0, 0, 0, 0), target("spirv.Image", void, 1, 1, 1, 0, 0, 0, 0)* %img7.addr, align 4
+  %call6 = call spir_func i32 @_Z24get_image_num_mip_levels26ocl_image2d_array_depth_ro(target("spirv.Image", void, 1, 1, 1, 0, 0, 0, 0) %6)
+  ret void
+}
+
+declare spir_func i32 @_Z24get_image_num_mip_levels14ocl_image1d_ro(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0))
+
+declare spir_func i32 @_Z24get_image_num_mip_levels14ocl_image2d_ro(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0))
+
+declare spir_func i32 @_Z24get_image_num_mip_levels14ocl_image3d_ro(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0))
+
+declare spir_func i32 @_Z24get_image_num_mip_levels20ocl_image1d_array_ro(target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 0))
+
+declare spir_func i32 @_Z24get_image_num_mip_levels20ocl_image2d_array_ro(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0))
+
+declare spir_func i32 @_Z24get_image_num_mip_levels20ocl_image2d_depth_ro(target("spirv.Image", void, 1, 1, 0, 0, 0, 0, 0))
+
+declare spir_func i32 @_Z24get_image_num_mip_levels26ocl_image2d_array_depth_ro(target("spirv.Image", void, 1, 1, 1, 0, 0, 0, 0))

llvm/test/CodeGen/SPIRV/types/image-query-size-with-access-qualifier.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; 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) {
+define spir_kernel void @sample_kernel(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0) %input) {
 entry:
-  %call = call spir_func i32 @_Z15get_image_width20ocl_image2d_array_ro(%opencl.image2d_array_ro_t addrspace(1)* %input)
+  %call = call spir_func i32 @_Z15get_image_width20ocl_image2d_array_ro(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0) %input)
   ret void
 }
 
-declare spir_func i32 @_Z15get_image_width20ocl_image2d_array_ro(%opencl.image2d_array_ro_t addrspace(1)*)
+declare spir_func i32 @_Z15get_image_width20ocl_image2d_array_ro(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0))

llvm/test/CodeGen/SPIRV/types/image-query-size.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; 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)
+define spir_kernel void @test_image1d(i32 addrspace(1)* nocapture %sizes, target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0) %img, target("spirv.Image", void, 5, 0, 0, 0, 0, 0, 0) %buffer, target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 0) %array) {
+  %1 = tail call spir_func i32 @_Z15get_image_width14ocl_image1d_ro(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0) %img)
+  %2 = tail call spir_func i32 @_Z15get_image_width21ocl_image1d_buffer_ro(target("spirv.Image", void, 5, 0, 0, 0, 0, 0, 0) %buffer)
+  %3 = tail call spir_func i32 @_Z15get_image_width20ocl_image1d_array_ro(target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 0) %array)
+  %4 = tail call spir_func i64 @_Z20get_image_array_size20ocl_image1d_array_ro(target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 0) %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_width14ocl_image1d_ro(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0))
 
-declare spir_func i32 @_Z15get_image_width21ocl_image1d_buffer_ro(%opencl.image1d_buffer_ro_t addrspace(1)*)
+declare spir_func i32 @_Z15get_image_width21ocl_image1d_buffer_ro(target("spirv.Image", void, 5, 0, 0, 0, 0, 0, 0))
 
-declare spir_func i32 @_Z15get_image_width20ocl_image1d_array_ro(%opencl.image1d_array_ro_t addrspace(1)*)
+declare spir_func i32 @_Z15get_image_width20ocl_image1d_array_ro(target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 0))
 
-declare spir_func i64 @_Z20get_image_array_size20ocl_image1d_array_ro(%opencl.image1d_array_ro_t addrspace(1)*)
+declare spir_func i64 @_Z20get_image_array_size20ocl_image1d_array_ro(target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 0))
 
-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)
+define spir_kernel void @test_image2d(i32 addrspace(1)* nocapture %sizes, target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %img, target("spirv.Image", void, 1, 1, 0, 0, 0, 0, 0) %img_depth, target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0) %array, target("spirv.Image", void, 1, 1, 1, 0, 0, 0, 0) %array_depth) {
+  %1 = tail call spir_func i32 @_Z15get_image_width14ocl_image2d_ro(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %img)
+  %2 = tail call spir_func i32 @_Z16get_image_height14ocl_image2d_ro(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %img)
+  %3 = tail call spir_func <2 x i32> @_Z13get_image_dim14ocl_image2d_ro(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %img)
+  %4 = tail call spir_func i32 @_Z15get_image_width20ocl_image2d_array_ro(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0) %array)
+  %5 = tail call spir_func i32 @_Z16get_image_height20ocl_image2d_array_ro(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0) %array)
+  %6 = tail call spir_func i64 @_Z20get_image_array_size20ocl_image2d_array_ro(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0) %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)
+  %8 = tail call spir_func <2 x i32> @_Z13get_image_dim20ocl_image2d_array_ro(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0) %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 @_Z15get_image_width14ocl_image2d_ro(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0))
 
-declare spir_func i32 @_Z16get_image_height14ocl_image2d_ro(%opencl.image2d_ro_t addrspace(1)*)
+declare spir_func i32 @_Z16get_image_height14ocl_image2d_ro(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0))
 
-declare spir_func <2 x i32> @_Z13get_image_dim14ocl_image2d_ro(%opencl.image2d_ro_t addrspace(1)*)
+declare spir_func <2 x i32> @_Z13get_image_dim14ocl_image2d_ro(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0))
 
-declare spir_func i32 @_Z15get_image_width20ocl_image2d_array_ro(%opencl.image2d_array_ro_t addrspace(1)*)
+declare spir_func i32 @_Z15get_image_width20ocl_image2d_array_ro(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0))
 
-declare spir_func i32 @_Z16get_image_height20ocl_image2d_array_ro(%opencl.image2d_array_ro_t addrspace(1)*)
+declare spir_func i32 @_Z16get_image_height20ocl_image2d_array_ro(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0))
 
-declare spir_func i64 @_Z20get_image_array_size20ocl_image2d_array_ro(%opencl.image2d_array_ro_t addrspace(1)*)
+declare spir_func i64 @_Z20get_image_array_size20ocl_image2d_array_ro(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0))
 
-declare spir_func <2 x i32> @_Z13get_image_dim20ocl_image2d_array_ro(%opencl.image2d_array_ro_t addrspace(1)*)
+declare spir_func <2 x i32> @_Z13get_image_dim20ocl_image2d_array_ro(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0))
 
-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)
+define spir_kernel void @test_image3d(i32 addrspace(1)* nocapture %sizes, target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %img) {
+  %1 = tail call spir_func i32 @_Z15get_image_width14ocl_image3d_ro(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %img)
+  %2 = tail call spir_func i32 @_Z16get_image_height14ocl_image3d_ro(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %img)
+  %3 = tail call spir_func i32 @_Z15get_image_depth14ocl_image3d_ro(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %img)
+  %4 = tail call spir_func <4 x i32> @_Z13get_image_dim14ocl_image3d_ro(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %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 @_Z15get_image_width14ocl_image3d_ro(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0))
 
-declare spir_func i32 @_Z16get_image_height14ocl_image3d_ro(%opencl.image3d_ro_t addrspace(1)*)
+declare spir_func i32 @_Z16get_image_height14ocl_image3d_ro(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0))
 
-declare spir_func i32 @_Z15get_image_depth14ocl_image3d_ro(%opencl.image3d_ro_t addrspace(1)*)
+declare spir_func i32 @_Z15get_image_depth14ocl_image3d_ro(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0))
 
-declare spir_func <4 x i32> @_Z13get_image_dim14ocl_image3d_ro(%opencl.image3d_ro_t addrspace(1)*)
+declare spir_func <4 x i32> @_Z13get_image_dim14ocl_image3d_ro(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0))
 
-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)
+define spir_kernel void @test_image2d_array_depth_t(i32 addrspace(1)* nocapture %sizes, target("spirv.Image", void, 1, 1, 1, 0, 0, 0, 0) %array) {
+  %1 = tail call spir_func i32 @_Z15get_image_width26ocl_image2d_array_depth_ro(target("spirv.Image", void, 1, 1, 1, 0, 0, 0, 0) %array)
+  %2 = tail call spir_func i32 @_Z16get_image_height26ocl_image2d_array_depth_ro(target("spirv.Image", void, 1, 1, 1, 0, 0, 0, 0) %array)
+  %3 = tail call spir_func i64 @_Z20get_image_array_size26ocl_image2d_array_depth_ro(target("spirv.Image", void, 1, 1, 1, 0, 0, 0, 0) %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 @_Z15get_image_width26ocl_image2d_array_depth_ro(target("spirv.Image", void, 1, 1, 1, 0, 0, 0, 0))
 
-declare spir_func i32 @_Z16get_image_height26ocl_image2d_array_depth_ro(%opencl.image2d_array_depth_ro_t addrspace(1)*)
+declare spir_func i32 @_Z16get_image_height26ocl_image2d_array_depth_ro(target("spirv.Image", void, 1, 1, 1, 0, 0, 0, 0))
 
-declare spir_func i64 @_Z20get_image_array_size26ocl_image2d_array_depth_ro(%opencl.image2d_array_depth_ro_t addrspace(1)*)
+declare spir_func i64 @_Z20get_image_array_size26ocl_image2d_array_depth_ro(target("spirv.Image", void, 1, 1, 1, 0, 0, 0, 0))
 No newline at end of file

llvm/test/CodeGen/SPIRV/types/image-read-ms.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; 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) {
+define spir_kernel void @sample_test(target("spirv.Image", void, 1, 0, 0, 1, 0, 0, 0) %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)
+  %call2 = tail call spir_func i32 @_Z15get_image_width19ocl_image2d_msaa_ro(target("spirv.Image", void, 1, 0, 0, 1, 0, 0, 0) %source)
+  %call3 = tail call spir_func i32 @_Z16get_image_height19ocl_image2d_msaa_ro(target("spirv.Image", void, 1, 0, 0, 1, 0, 0, 0) %source)
+  %call4 = tail call spir_func i32 @_Z21get_image_num_samples19ocl_image2d_msaa_ro(target("spirv.Image", void, 1, 0, 0, 1, 0, 0, 0) %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)
+  %call9 = tail call spir_func <4 x float> @_Z11read_imagef19ocl_image2d_msaa_roDv2_ii(target("spirv.Image", void, 1, 0, 0, 1, 0, 0, 0) %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 @_Z15get_image_width19ocl_image2d_msaa_ro(target("spirv.Image", void, 1, 0, 0, 1, 0, 0, 0))
 
-declare spir_func i32 @_Z16get_image_height19ocl_image2d_msaa_ro(%opencl.image2d_msaa_ro_t addrspace(1)*)
+declare spir_func i32 @_Z16get_image_height19ocl_image2d_msaa_ro(target("spirv.Image", void, 1, 0, 0, 1, 0, 0, 0))
 
-declare spir_func i32 @_Z21get_image_num_samples19ocl_image2d_msaa_ro(%opencl.image2d_msaa_ro_t addrspace(1)*)
+declare spir_func i32 @_Z21get_image_num_samples19ocl_image2d_msaa_ro(target("spirv.Image", void, 1, 0, 0, 1, 0, 0, 0))
 
-declare spir_func <4 x float> @_Z11read_imagef19ocl_image2d_msaa_roDv2_ii(%opencl.image2d_msaa_ro_t addrspace(1)*, <2 x i32>, i32)
+declare spir_func <4 x float> @_Z11read_imagef19ocl_image2d_msaa_roDv2_ii(target("spirv.Image", void, 1, 0, 0, 1, 0, 0, 0), <2 x i32>, i32)

llvm/test/CodeGen/SPIRV/types/image-read.ll

+; 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));
+;; }
+
+define dso_local spir_kernel void @kernelA(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %input) {
+entry:
+  %input.addr = alloca target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0), align 8
+  %c = alloca <4 x i32>, align 16
+  %.compoundliteral = alloca <4 x i32>, align 16
+  store target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %input, ptr %input.addr, align 8
+  %0 = load target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0), ptr %input.addr, align 8
+  store <4 x i32> zeroinitializer, ptr %.compoundliteral, align 16
+  %1 = load <4 x i32>, ptr %.compoundliteral, align 16
+  %call = call spir_func <4 x i32> @_Z12read_imageui14ocl_image3d_roDv4_i(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %0, <4 x i32> noundef %1)
+  store <4 x i32> %call, ptr %c, align 16
+  ret void
+}
+
+declare spir_func <4 x i32> @_Z12read_imageui14ocl_image3d_roDv4_i(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %0, <4 x i32> noundef %1)
+
+define dso_local spir_kernel void @kernelB(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %input) {
+entry:
+  %input.addr = alloca target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0), align 8
+  %f = alloca <4 x float>, align 16
+  %.compoundliteral = alloca <4 x i32>, align 16
+  store target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %input, ptr %input.addr, align 8
+  %0 = load target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0), ptr %input.addr, align 8
+  store <4 x i32> zeroinitializer, ptr %.compoundliteral, align 16
+  %1 = load <4 x i32>, ptr %.compoundliteral, align 16
+  %call = call spir_func <4 x float> @_Z11read_imagef14ocl_image3d_roDv4_i(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %0, <4 x i32> noundef %1)
+  store <4 x float> %call, ptr %f, align 16
+  ret void
+}
+
+declare spir_func <4 x float> @_Z11read_imagef14ocl_image3d_roDv4_i(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %0, <4 x i32> noundef %1)

llvm/test/CodeGen/SPIRV/types/image-sample-explicit-level-of-details.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; 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) {
+define spir_kernel void @sample_kernel(target("spirv.Image", void, 1, 1, 0, 0, 0, 0, 0) %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.tmp1 = call spir_func <2 x i32> @_Z13get_image_dim20ocl_image2d_depth_ro(target("spirv.Image", void, 1, 1, 0, 0, 0, 0, 0) %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)
+  %call6.tmp.tmp = call spir_func float @_Z11read_imagef20ocl_image2d_depth_ro11ocl_samplerDv2_i(target("spirv.Image", void, 1, 1, 0, 0, 0, 0, 0) %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 float @_Z11read_imagef20ocl_image2d_depth_ro11ocl_samplerDv2_i(target("spirv.Image", void, 1, 1, 0, 0, 0, 0, 0), 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)*)
+declare spir_func <2 x i32> @_Z13get_image_dim20ocl_image2d_depth_ro(target("spirv.Image", void, 1, 1, 0, 0, 0, 0, 0))
 No newline at end of file

llvm/test/CodeGen/SPIRV/types/image-sampled.ll

+; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+
+; CHECK: %[[#image1d_t:]] = OpTypeImage
+; CHECK: %[[#sampler_t:]] = OpTypeSampler
+; CHECK: %[[#sampled_image_t:]] = OpTypeSampledImage
+
+declare dso_local spir_func ptr addrspace(4) @_Z20__spirv_SampledImageI14ocl_image1d_roPvET0_T_11ocl_sampler(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0) %0, target("spirv.Sampler") %1) local_unnamed_addr
+
+declare dso_local spir_func <4 x float> @_Z30__spirv_ImageSampleExplicitLodIPvDv4_fiET0_T_T1_if(ptr addrspace(4) %0, i32 %1, i32 %2, float %3) local_unnamed_addr
+
+@__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(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0), target("spirv.Sampler")) {
+; CHECK: OpFunction
+; CHECK: %[[#image:]] = OpFunctionParameter %[[#image1d_t]]
+; CHECK: %[[#sampler:]] = OpFunctionParameter %[[#sampler_t]]
+  %3 = load <3 x i64>, ptr addrspace(2) @__spirv_BuiltInGlobalInvocationId, align 32
+  %4 = extractelement <3 x i64> %3, i64 0
+  %5 = trunc i64 %4 to i32
+  %6 = call spir_func ptr addrspace(4) @_Z20__spirv_SampledImageI14ocl_image1d_roPvET0_T_11ocl_sampler(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0) %0, target("spirv.Sampler") %1)
+  %7 = call spir_func <4 x float> @_Z30__spirv_ImageSampleExplicitLodIPvDv4_fiET0_T_T1_if(ptr 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
+}

llvm/test/CodeGen/SPIRV/types/image-sampled2.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; 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 {
+define dso_local spir_kernel void @sample_kernel_float(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %input, <2 x float> noundef %coords, <4 x float> addrspace(1)* nocapture noundef writeonly %results, target("spirv.Sampler") %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)
+  %0 = tail call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 32)
+  %call = tail call spir_func <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv2_f(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %input, target("spirv.Sampler") %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)
+  %call1 = tail call spir_func <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv2_f(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %input, target("spirv.Sampler") %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)
+  %1 = tail call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 22)
+  %call2 = tail call spir_func <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv2_f(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %input, target("spirv.Sampler") %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 <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv2_f(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0), target("spirv.Sampler"), <2 x float> noundef) local_unnamed_addr
 
-declare spir_func %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32) local_unnamed_addr
+declare spir_func target("spirv.Sampler") @__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 {
+define dso_local spir_kernel void @sample_kernel_int(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %input, <2 x float> noundef %coords, <4 x i32> addrspace(1)* nocapture noundef writeonly %results, target("spirv.Sampler") %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)
+  %0 = tail call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 32)
+  %call = tail call spir_func <4 x i32> @_Z11read_imagei14ocl_image2d_ro11ocl_samplerDv2_f(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %input, target("spirv.Sampler") %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)
+  %call1 = tail call spir_func <4 x i32> @_Z11read_imagei14ocl_image2d_ro11ocl_samplerDv2_f(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %input, target("spirv.Sampler") %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)
+  %1 = tail call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 22)
+  %call2 = tail call spir_func <4 x i32> @_Z11read_imagei14ocl_image2d_ro11ocl_samplerDv2_f(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %input, target("spirv.Sampler") %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
+declare spir_func <4 x i32> @_Z11read_imagei14ocl_image2d_ro11ocl_samplerDv2_f(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0), target("spirv.Sampler"), <2 x float> noundef) local_unnamed_addr

llvm/test/CodeGen/SPIRV/types/image-unoptimized.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
+; 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]]

 ;;
 ;;
 ;; __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) {
+define dso_local spir_kernel void @test_fn(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %srcimg, target("spirv.Sampler") %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
+  %srcimg.addr = alloca target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0), align 4
+  %sampler.addr = alloca target("spirv.Sampler"), 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 target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %srcimg, target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0)* %srcimg.addr, align 4
+  store target("spirv.Sampler") %sampler, target("spirv.Sampler")* %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
+  %0 = load target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0), target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0)* %srcimg.addr, align 4
+  %1 = load target("spirv.Sampler"), target("spirv.Sampler")* %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)
+  %call3 = call spir_func <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv2_i(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %0, target("spirv.Sampler") %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)
+  %8 = load target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0), target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0)* %srcimg.addr, align 4
+  %call4 = call spir_func i32 @_Z15get_image_width14ocl_image2d_ro(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %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 <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv2_i(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0), target("spirv.Sampler"), <2 x i32> noundef)
 
-declare spir_func i32 @_Z15get_image_width14ocl_image2d_ro(%opencl.image2d_ro_t addrspace(1)*)
+declare spir_func i32 @_Z15get_image_width14ocl_image2d_ro(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0))

llvm/test/CodeGen/SPIRV/types/image-various-dimensions-access-qualifiers.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
 
 ;; FIXME: Write tests to ensure invalid usage of image are rejected, such as:
 ;;  - invalid AS (only global is allowed);
 ;;  - used in struct, union, array, pointer or return types;
 ;;  - used with invalid CV-qualifiers (const or volatile in C99).
 ;; FIXME: Write further tests to cover _array, _buffer, _depth, ... types.
 
-%opencl.image1d_ro_t = type opaque ;; read_only image1d_t
-%opencl.image2d_wo_t = type opaque ;; write_only image2d_t
-%opencl.image3d_rw_t = type opaque ;; read_write image3d_t
-
 define void @foo(
-  %opencl.image1d_ro_t addrspace(1)* %a,
-  %opencl.image2d_wo_t addrspace(1)* %b,
-  %opencl.image3d_rw_t addrspace(1)* %c,
+  target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0) %a,
+  target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 1) %b,
+  target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 2) %c,
   i32 addrspace(1)* %d
 ) {
-  %pixel = call <4 x i32> @_Z11read_imagei14ocl_image1d_roi(%opencl.image1d_ro_t addrspace(1)* %a, i32 0)
-  call void @_Z12write_imagei14ocl_image2d_woDv2_iDv4_i(%opencl.image2d_wo_t addrspace(1)* %b, <2 x i32> zeroinitializer, <4 x i32> %pixel)
-  %size = call i32 @_Z15get_image_width14ocl_image3d_rw(%opencl.image3d_rw_t addrspace(1)* %c)
+  %pixel = call <4 x i32> @_Z11read_imagei14ocl_image1d_roi(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0) %a, i32 0)
+  call void @_Z12write_imagei14ocl_image2d_woDv2_iDv4_i(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 1) %b, <2 x i32> zeroinitializer, <4 x i32> %pixel)
+  %size = call i32 @_Z15get_image_width14ocl_image3d_rw(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 2) %c)
   store i32 %size, i32 addrspace(1)* %d
   ret void
 }
 
-declare <4 x i32> @_Z11read_imagei14ocl_image1d_roi(%opencl.image1d_ro_t addrspace(1)*, i32)
+declare <4 x i32> @_Z11read_imagei14ocl_image1d_roi(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0), i32)
 
-declare void @_Z12write_imagei14ocl_image2d_woDv2_iDv4_i(%opencl.image2d_wo_t addrspace(1)*, <2 x i32>, <4 x i32>)
+declare void @_Z12write_imagei14ocl_image2d_woDv2_iDv4_i(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 1), <2 x i32>, <4 x i32>)
 
-declare i32 @_Z15get_image_width14ocl_image3d_rw(%opencl.image3d_rw_t addrspace(1)*)
+declare i32 @_Z15get_image_width14ocl_image3d_rw(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 2))
 
 
 ;; Capabilities:
 ; CHECK-DAG: OpCapability ImageReadWrite
 ; CHECK-NOT: DAG-FENCE
 
 ;; Types, Constants and Variables:
 ;; FIXME: The values should be double checked here.

llvm/test/CodeGen/SPIRV/types/image-write.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; 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 {
+define dso_local spir_kernel void @test_img2d(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 1) %image_wo, target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 2) %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)
+  call spir_func void @_Z12write_imagef14ocl_image2d_woDv2_iDv4_f(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 1) %image_wo, <2 x i32> noundef zeroinitializer, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei14ocl_image2d_woDv2_iDv4_i(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 1) %image_wo, <2 x i32> noundef zeroinitializer, <4 x i32> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagef14ocl_image2d_rwDv2_iDv4_f(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 2) %image_rw, <2 x i32> noundef zeroinitializer, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei14ocl_image2d_rwDv2_iDv4_i(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 2) %image_rw, <2 x i32> noundef zeroinitializer, <4 x i32> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagef14ocl_image2d_woDv2_iiDv4_f(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 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(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 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_imagef14ocl_image2d_woDv2_iDv4_f(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 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_imagei14ocl_image2d_woDv2_iDv4_i(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 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_imagef14ocl_image2d_rwDv2_iDv4_f(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 2), <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_imagei14ocl_image2d_rwDv2_iDv4_i(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 2), <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_imagef14ocl_image2d_woDv2_iiDv4_f(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 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
+declare spir_func void @_Z12write_imagei14ocl_image2d_woDv2_iiDv4_i(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 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 {
+define dso_local spir_kernel void @test_img2d_array(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 1) %image_wo, target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 2) %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)
+  call spir_func void @_Z12write_imagef20ocl_image2d_array_woDv4_iDv4_f(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 1) %image_wo, <4 x i32> noundef zeroinitializer, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei20ocl_image2d_array_woDv4_iS0_(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 1) %image_wo, <4 x i32> noundef zeroinitializer, <4 x i32> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagef20ocl_image2d_array_rwDv4_iDv4_f(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 2) %image_rw, <4 x i32> noundef zeroinitializer, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei20ocl_image2d_array_rwDv4_iS0_(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 2) %image_rw, <4 x i32> noundef zeroinitializer, <4 x i32> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagef20ocl_image2d_array_woDv4_iiDv4_f(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 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_(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 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_imagef20ocl_image2d_array_woDv4_iDv4_f(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 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_imagei20ocl_image2d_array_woDv4_iS0_(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 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_imagef20ocl_image2d_array_rwDv4_iDv4_f(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 2), <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_imagei20ocl_image2d_array_rwDv4_iS0_(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 2), <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_imagef20ocl_image2d_array_woDv4_iiDv4_f(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 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
+declare spir_func void @_Z12write_imagei20ocl_image2d_array_woDv4_iiS0_(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 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 {
+define dso_local spir_kernel void @test_img1d(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 1) %image_wo, target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 2) %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)
+  call spir_func void @_Z12write_imagef14ocl_image1d_woiDv4_f(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 1) %image_wo, i32 noundef 0, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei14ocl_image1d_woiDv4_i(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 1) %image_wo, i32 noundef 0, <4 x i32> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagef14ocl_image1d_rwiDv4_f(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 2) %image_rw, i32 noundef 0, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei14ocl_image1d_rwiDv4_i(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 2) %image_rw, i32 noundef 0, <4 x i32> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagef14ocl_image1d_woiiDv4_f(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 1) %image_wo, i32 noundef 0, i32 noundef 0, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei14ocl_image1d_woiiDv4_i(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 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_imagef14ocl_image1d_woiDv4_f(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 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_imagei14ocl_image1d_woiDv4_i(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 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_imagef14ocl_image1d_rwiDv4_f(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 2), 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_imagei14ocl_image1d_rwiDv4_i(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 2), 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_imagef14ocl_image1d_woiiDv4_f(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 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
+declare spir_func void @_Z12write_imagei14ocl_image1d_woiiDv4_i(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 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 {
+define dso_local spir_kernel void @test_img1d_buffer(target("spirv.Image", void, 5, 0, 0, 0, 0, 0, 1) %image_wo, target("spirv.Image", void, 5, 0, 0, 0, 0, 0, 2) %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)
+  call spir_func void @_Z12write_imagef21ocl_image1d_buffer_woiDv4_f(target("spirv.Image", void, 5, 0, 0, 0, 0, 0, 1) %image_wo, i32 noundef 0, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei21ocl_image1d_buffer_woiDv4_i(target("spirv.Image", void, 5, 0, 0, 0, 0, 0, 1) %image_wo, i32 noundef 0, <4 x i32> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagef21ocl_image1d_buffer_rwiDv4_f(target("spirv.Image", void, 5, 0, 0, 0, 0, 0, 2) %image_rw, i32 noundef 0, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei21ocl_image1d_buffer_rwiDv4_i(target("spirv.Image", void, 5, 0, 0, 0, 0, 0, 2) %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_imagef21ocl_image1d_buffer_woiDv4_f(target("spirv.Image", void, 5, 0, 0, 0, 0, 0, 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_imagei21ocl_image1d_buffer_woiDv4_i(target("spirv.Image", void, 5, 0, 0, 0, 0, 0, 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_imagef21ocl_image1d_buffer_rwiDv4_f(target("spirv.Image", void, 5, 0, 0, 0, 0, 0, 2), 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
+declare spir_func void @_Z12write_imagei21ocl_image1d_buffer_rwiDv4_i(target("spirv.Image", void, 5, 0, 0, 0, 0, 0, 2), 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 {
+define dso_local spir_kernel void @test_img1d_array(target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 1) %image_wo, target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 2) %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)
+  call spir_func void @_Z12write_imagef20ocl_image1d_array_woDv2_iDv4_f(target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 1) %image_wo, <2 x i32> noundef zeroinitializer, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei20ocl_image1d_array_woDv2_iDv4_i(target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 1) %image_wo, <2 x i32> noundef zeroinitializer, <4 x i32> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagef20ocl_image1d_array_rwDv2_iDv4_f(target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 2) %image_rw, <2 x i32> noundef zeroinitializer, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei20ocl_image1d_array_rwDv2_iDv4_i(target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 2) %image_rw, <2 x i32> noundef zeroinitializer, <4 x i32> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagef20ocl_image1d_array_woDv2_iiDv4_f(target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 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(target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 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_imagef20ocl_image1d_array_woDv2_iDv4_f(target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 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_imagei20ocl_image1d_array_woDv2_iDv4_i(target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 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_imagef20ocl_image1d_array_rwDv2_iDv4_f(target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 2), <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_imagei20ocl_image1d_array_rwDv2_iDv4_i(target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 2), <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_imagef20ocl_image1d_array_woDv2_iiDv4_f(target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 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
+declare spir_func void @_Z12write_imagei20ocl_image1d_array_woDv2_iiDv4_i(target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 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 {
+define dso_local spir_kernel void @test_img2d_depth(target("spirv.Image", void, 1, 1, 0, 0, 0, 0, 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)
+  call spir_func void @_Z12write_imagef20ocl_image2d_depth_woDv2_if(target("spirv.Image", void, 1, 1, 0, 0, 0, 0, 1) %image_wo, <2 x i32> noundef zeroinitializer, float noundef 0.000000e+00)
+  call spir_func void @_Z12write_imagef20ocl_image2d_depth_woDv2_if(target("spirv.Image", void, 1, 1, 0, 0, 0, 0, 1) %image_wo, <2 x i32> noundef zeroinitializer, float noundef 0.000000e+00)
+  call spir_func void @_Z12write_imagef20ocl_image2d_depth_woDv2_iif(target("spirv.Image", void, 1, 1, 0, 0, 0, 0, 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_if(target("spirv.Image", void, 1, 1, 0, 0, 0, 0, 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
+declare spir_func void @_Z12write_imagef20ocl_image2d_depth_woDv2_iif(target("spirv.Image", void, 1, 1, 0, 0, 0, 0, 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 {
+define dso_local spir_kernel void @test_img2d_array_depth(target("spirv.Image", void, 1, 1, 1, 0, 0, 0, 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)
+  call spir_func void @_Z12write_imagef26ocl_image2d_array_depth_woDv4_if(target("spirv.Image", void, 1, 1, 1, 0, 0, 0, 1) %image_wo, <4 x i32> noundef zeroinitializer, float noundef 0.000000e+00)
+  call spir_func void @_Z12write_imagef26ocl_image2d_array_depth_woDv4_if(target("spirv.Image", void, 1, 1, 1, 0, 0, 0, 1) %image_wo, <4 x i32> noundef zeroinitializer, float noundef 0.000000e+00)
+  call spir_func void @_Z12write_imagef26ocl_image2d_array_depth_woDv4_iif(target("spirv.Image", void, 1, 1, 1, 0, 0, 0, 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_if(target("spirv.Image", void, 1, 1, 1, 0, 0, 0, 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
+declare spir_func void @_Z12write_imagef26ocl_image2d_array_depth_woDv4_iif(target("spirv.Image", void, 1, 1, 1, 0, 0, 0, 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 {
+define dso_local spir_kernel void @test_img3d(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 1) %image_wo, target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 2) %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)
+  call spir_func void @_Z12write_imagef14ocl_image3d_woDv4_iDv4_f(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 1) %image_wo, <4 x i32> noundef zeroinitializer, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei14ocl_image3d_woDv4_iS0_(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 1) %image_wo, <4 x i32> noundef zeroinitializer, <4 x i32> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagef14ocl_image3d_rwDv4_iDv4_f(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 2) %image_rw, <4 x i32> noundef zeroinitializer, <4 x float> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagei14ocl_image3d_rwDv4_iS0_(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 2) %image_rw, <4 x i32> noundef zeroinitializer, <4 x i32> noundef zeroinitializer)
+  call spir_func void @_Z12write_imagef14ocl_image3d_woDv4_iiDv4_f(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 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_(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 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_imagef14ocl_image3d_woDv4_iDv4_f(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 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_imagei14ocl_image3d_woDv4_iS0_(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 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_imagef14ocl_image3d_rwDv4_iDv4_f(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 2), <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_imagei14ocl_image3d_rwDv4_iS0_(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 2), <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_imagef14ocl_image3d_woDv4_iiDv4_f(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 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
+declare spir_func void @_Z12write_imagei14ocl_image3d_woDv4_iiS0_(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 1), <4 x i32> noundef, i32 noundef, <4 x i32> noundef) local_unnamed_addr

llvm/test/CodeGen/SPIRV/types/image.ll

-; RUN: llc -O0 -opaque-pointers=0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; 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:
+define spir_kernel void @image_copy(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %image1, target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %image2) !kernel_arg_access_qual !1 {
   %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)
+  %call4 = tail call spir_func <4 x float> @_Z11read_imagef11ocl_image2d11ocl_samplerDv2_i(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %image1, i32 20, <2 x i32> %vecinit3)
+  tail call spir_func void @_Z12write_imagef11ocl_image2dDv2_iDv4_f(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %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 <4 x float> @_Z11read_imagef11ocl_image2d11ocl_samplerDv2_i(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0), i32, <2 x i32>)
 
-declare spir_func void @_Z12write_imagef11ocl_image2dDv2_iDv4_f(%opencl.image2d_t addrspace(1)*, <2 x i32>, <4 x float>)
+declare spir_func void @_Z12write_imagef11ocl_image2dDv2_iDv4_f(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0), <2 x i32>, <4 x float>)
 
 !1 = !{!"read_only", !"write_only"}

llvm/test/CodeGen/SPIRV/types/image1d.ll

+; 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
+
+define spir_kernel void @test_image_dim(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0) %image1d, target("spirv.Image", void, 5, 0, 0, 0, 0, 0, 0) %image1d_buffer) {
+  ret void
+}

llvm/test/CodeGen/SPIRV/types/pipe-sampler-image.ll

 ;;  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 -opaque-pointers=0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; 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) {
+  target("spirv.Pipe", 0) %a,
+  target("spirv.Pipe", 1) %b,
+  target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0) %c1,
+  target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %d1,
+  target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %e1,
+  target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0) %f1,
+  target("spirv.Image", void, 5, 0, 0, 0, 0, 0, 0) %g1,
+  target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 1) %c2,
+  target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 2) %d3,
+  target("spirv.Sampler") %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)
+  %.tmp = call spir_func <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv4_if(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %d1, target("spirv.Sampler") %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)
+  %0 = call target("spirv.Sampler") @__translate_sampler_initializer(i32 32)
+  %.tmp2 = call spir_func <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv4_if(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %d1, target("spirv.Sampler") %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 spir_func <4 x float> @_Z11read_imagef14ocl_image2d_ro11ocl_samplerDv4_if(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0), target("spirv.Sampler"), <4 x i32>, float)
 
-declare %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32)
+declare target("spirv.Sampler") @__translate_sampler_initializer(i32)