[SPIR-V] [WIP] Convert tests to opaque pointers

[mpaszkowski]

Still WIP, but CTS basic/* tests in both opaque pointer mode and typed pointer mode are now fixed. We still have failures in CTS printf/* tests.

[mpaszkowski]

New final revision can be found here: https://reviews.llvm.org/D156049

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 {
+using SPIRVType = const MachineInstr;
+
 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 {

 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.h

SPIRVType *getOrCreateSPIRVVectorType(SPIRVType *BaseType,
                                         MachineIRBuilder &MIRBuilder);
   SPIRVType *getOrCreateSPIRVVectorType(SPIRVType *BaseType,
                                         unsigned NumElements, MachineInstr &I,
                                         const SPIRVInstrInfo &TII);
   SPIRVType *getOrCreateSPIRVArrayType(SPIRVType *BaseType,
                                        unsigned NumElements, MachineInstr &I,
                                        const SPIRVInstrInfo &TII);
 
+  // TODO: Remove once typed pointers are no longer supported
   SPIRVType *getOrCreateSPIRVPointerType(
       SPIRVType *BaseType, MachineIRBuilder &MIRBuilder,
       SPIRV::StorageClass::StorageClass SClass = SPIRV::StorageClass::Function);
+
   SPIRVType *getOrCreateSPIRVPointerType(
       SPIRVType *BaseType, MachineInstr &I, const SPIRVInstrInfo &TII,
       SPIRV::StorageClass::StorageClass SClass = SPIRV::StorageClass::Function);
 
   SPIRVType *
   getOrCreateOpTypeImage(MachineIRBuilder &MIRBuilder, SPIRVType *SampledType,
                          SPIRV::Dim::Dim Dim, uint32_t Depth, uint32_t Arrayed,
                          uint32_t Multisampled, uint32_t Sampled,

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 if (Ty->isOpaquePointerTy())
+    Reg =
+        DT.find(Type::getInt8Ty(MIRBuilder.getMF().getFunction().getContext()),
+                Ty->getPointerAddressSpace(), &MIRBuilder.getMF());
+  else
+    Reg = DT.find(Ty->getNonOpaquePointerElementType(),
+                  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);
 }
 
 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();
+        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/SampledImageRetType.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
 
-%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 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(i8 addrspace(4)*, i32, i32, float) 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(%opencl.image1d_ro_t addrspace(1)*, %opencl.sampler_t addrspace(2)*) {
+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>, <3 x i64> addrspace(2)* @__spirv_BuiltInGlobalInvocationId, align 32
+  %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 = 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)
+  %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/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/constant/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/event_no_group_cap.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/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/image_decl_func_arg.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:     %[[#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:
+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(%opencl.image2d_ro_t addrspace(1)* %v2, <2 x float> %v3) {
-entry:
+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(%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>)
+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/image_dim.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: 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:
+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/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/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/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/multi_md.ll

 ;; Check duplicate operands in opencl.ocl.version metadata is accepted without
 ;; assertion.
 
 ; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
 
 %struct.my_struct_t = type { i8, i32 }
 
 @var = addrspace(1) global %struct.my_struct_t { i8 97, i32 42 }, align 4
 
 define spir_kernel void @__OpenCL_writer_kernel(i8 zeroext %c, i32 %i) {
 entry:
   %c.addr = alloca i8, align 1
   %i.addr = alloca i32, align 4
-  store i8 %c, i8* %c.addr, align 1
-  store i32 %i, i32* %i.addr, align 4
-  %0 = load i8, i8* %c.addr, align 1
-  store i8 %0, i8 addrspace(1)* getelementptr inbounds (%struct.my_struct_t, %struct.my_struct_t addrspace(1)* @var, i32 0, i32 0), align 1
-  %1 = load i32, i32* %i.addr, align 4
-  store i32 %1, i32 addrspace(1)* getelementptr inbounds (%struct.my_struct_t, %struct.my_struct_t addrspace(1)* @var, i32 0, i32 1), align 4
+  store i8 %c, ptr %c.addr, align 1
+  store i32 %i, ptr %i.addr, align 4
+  %0 = load i8, ptr %c.addr, align 1
+  store i8 %0, ptr addrspace(1) @var, align 1
+  %1 = load i32, ptr %i.addr, align 4
+  %2 = getelementptr inbounds %struct.my_struct_t, ptr addrspace(1) @var, i32 0, i32 1
+  store i32 %1, ptr addrspace(1) %2, align 4
   ret void
 }
 
-define spir_kernel void @__OpenCL_reader_kernel(i8 addrspace(1)* %C, i32 addrspace(1)* %I) {
+define spir_kernel void @__OpenCL_reader_kernel(ptr addrspace(1) %C, ptr addrspace(1) %I) {
 entry:
-  %C.addr = alloca i8 addrspace(1)*, align 8
-  %I.addr = alloca i32 addrspace(1)*, align 8
-  store i8 addrspace(1)* %C, i8 addrspace(1)** %C.addr, align 8
-  store i32 addrspace(1)* %I, i32 addrspace(1)** %I.addr, align 8
-  %0 = load i8, i8 addrspace(1)* getelementptr inbounds (%struct.my_struct_t, %struct.my_struct_t addrspace(1)* @var, i32 0, i32 0), align 1
-  %1 = load i8 addrspace(1)*, i8 addrspace(1)** %C.addr, align 8
-  store i8 %0, i8 addrspace(1)* %1, align 1
-  %2 = load i32, i32 addrspace(1)* getelementptr inbounds (%struct.my_struct_t, %struct.my_struct_t addrspace(1)* @var, i32 0, i32 1), align 4
-  %3 = load i32 addrspace(1)*, i32 addrspace(1)** %I.addr, align 8
-  store i32 %2, i32 addrspace(1)* %3, align 4
+  %C.addr = alloca ptr addrspace(1), align 8
+  %I.addr = alloca ptr addrspace(1), align 8
+  store ptr addrspace(1) %C, ptr %C.addr, align 8
+  store ptr addrspace(1) %I, ptr %I.addr, align 8
+  %0 = load i8, ptr addrspace(1) @var, align 1
+  %1 = load ptr addrspace(1), ptr %C.addr, align 8
+  store i8 %0, ptr addrspace(1) %1, align 1
+  %2 = getelementptr inbounds %struct.my_struct_t, ptr addrspace(1) @var, i32 0, i32 1
+  %3 = load i32, ptr addrspace(1) %2, align 4
+  %4 = load ptr addrspace(1), ptr %I.addr, align 8
+  store i32 %3, ptr addrspace(1) %4, align 4
   ret void
 }
 
 ;; "cl_images" should be encoded as BasicImage capability,
 ;; but images are not used in this test case, so this capability is not required.
 ; CHECK-NOT: OpExtension "cl_images"
 ; CHECK-DAG: OpSourceExtension "cl_khr_int64_base_atomics"
 ; CHECK-DAG: OpSourceExtension "cl_khr_int64_extended_atomics"
 ; CHECK:     OpSource OpenCL_C 200000
 
 !opencl.ocl.version = !{!13, !13}
 !opencl.used.extensions = !{!24, !25}
 
 !13 = !{i32 2, i32 0}
 !24 = !{!"cl_khr_int64_base_atomics"}
 !25 = !{!"cl_khr_int64_base_atomics", !"cl_khr_int64_extended_atomics"}

llvm/test/CodeGen/SPIRV/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/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/opencl/basic/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/read_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: %[[#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) {
+define dso_local spir_kernel void @kernelA(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %input) {
 entry:
-  %input.addr = alloca %opencl.image3d_ro_t addrspace(1)*, align 8
+  %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 %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
+  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(%opencl.image3d_ro_t addrspace(1)*, <4 x i32> noundef)
+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(%opencl.image3d_ro_t addrspace(1)* %input) {
+define dso_local spir_kernel void @kernelB(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0) %input) {
 entry:
-  %input.addr = alloca %opencl.image3d_ro_t addrspace(1)*, align 8
+  %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 %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
+  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(%opencl.image3d_ro_t addrspace(1)*, <4 x i32> noundef)
+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/transcoding/OpImageQuerySize.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))

llvm/test/CodeGen/SPIRV/transcoding/OpImageSampleExplicitLod.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))

llvm/test/CodeGen/SPIRV/transcoding/OpImageWrite.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/transcoding/SampledImage.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/transcoding/check_ro_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: %[[#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/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);
 ;;   constant block_t b1 = ^(int i) { return i + 1; };
 ;;   *res = b1(5);
 ;; }
 
-; CHECK-SPIRV1_4:   OpEntryPoint Kernel %[[#]] "block_kernel" %[[#InterfaceId:]]
-; CHECK-SPIRV1_4:   OpName %[[#InterfaceId]] "__block_literal_global"
+; CHECK-SPIRV1_4:   OpEntryPoint Kernel %[[#]] "block_kernel"
+; CHECK-SPIRV1_4:   OpName %[[#InterfaceId:]] "__block_literal_global"
 ; CHECK-SPIRV:      OpName %[[#block_invoke:]] "_block_invoke"
-; CHECK-SPIRV:      %[[#int:]] = OpTypeInt 32
 ; CHECK-SPIRV:      %[[#int8:]] = OpTypeInt 8
+; CHECK-SPIRV:      %[[#int:]] = OpTypeInt 32
 ; CHECK-SPIRV:      %[[#int8Ptr:]] = OpTypePointer Generic %[[#int8]]
 ; CHECK-SPIRV:      %[[#block_invoke_type:]] = OpTypeFunction %[[#int]] %[[#int8Ptr]] %[[#int]]
 ; CHECK-SPIRV:      %[[#five:]] = OpConstant %[[#int]] 5
 
 ; CHECK-SPIRV:      %[[#]] = OpFunctionCall %[[#int]] %[[#block_invoke]] %[[#]] %[[#five]]
 
 ; CHECK-SPIRV:      %[[#block_invoke]] = OpFunction %[[#int]] DontInline %[[#block_invoke_type]]
 ; CHECK-SPIRV-NEXT: %[[#]] = OpFunctionParameter %[[#int8Ptr]]

llvm/test/CodeGen/SPIRV/transcoding/optional-core-features-multiple.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/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: %[[#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-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:     %[[#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: %[[#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)