[Clang][SPIR-V] Emit target extension types for OpenCL types on SPIR-V.

[jcranmer-intel]

Most of the testing for this change has been in conjunction with the changes in the SPIRV-LLVM-Translator repository here: https://github.com/KhronosGroup/SPIRV-LLVM-Translator/pull/1799. I haven't updated the in-tree experimental target to support the target-extension types, but I did want to provide more documentation as to how these types work in the existing LLVM documentation for the SPIR-V backend.

There's still a few issues with this patch. First, I've made the use of these types unconditional for the SPIR-V target--it may be more appropriate to make these triggered off of a hidden option defaulted to off for now, or maybe based on whether or not opaque pointers are enabled. The other issue is the cast_image.cl test which... does an illegal operation in SPIR-V, and looking at the history of the test, doesn't seem to really test what it's claiming to test (https://reviews.llvm.org/D22927 was the revision it was adding tests for). I disabled the test for now for SPIR-V out of ignorance of which way to go about for fixing this test.

[arsenm]

clang/test/CodeGenOpenCL/cast_image.cl

2

Broken run line is not the way. should have explicit check for the failure

[bader]

@jcranmer-intel, thanks a lot for working on this. I'm so excited to see these changes!
Overall, it looks good to me, but I'd like to avoid some runtime computations if possible.

clang/lib/CodeGen/CGOpenCLRuntime.cpp

40

I believe this can be done at "compile time" (i.e. during the clang build, not clang run).
Can we have a pre-computed map from an OpenCL built-in type to a SPIR-V type?
Another option is compile-time evaluated function. This should be possible, right?

If I get it right, here we take a string representation of an OpenCL image type and process it at runtime, which seems to be unnecessary as we have pre-defined (by the spec) set of the types.

[jcranmer-intel]

clang/lib/CodeGen/CGOpenCLRuntime.cpp

40

I can definitely switch the read suffix to use a compile-time enum, since there are only 3 cases (plus, it's a static assert). Making the openCL name to int param conversion be a compile-time constant might be doable with some tricks, but I'll have to think about it for a little bit. It's a little harder because we're taking a string to 6 array values.

[bader]

clang/lib/CodeGen/CGOpenCLRuntime.cpp

40

I was going to suggest ripping of https://reviews.llvm.org/D108034, but it looks like it produces types which have OpenCL names with __spirv_* prefix. So unfortunately, I don't have a good example. The only thing coming to my mind is to build another table with SPIR-V type names, which can be obtained via OpenCL type id (offset?).

[svenvh]

it may be more appropriate to make these triggered off of a hidden option defaulted to off for now, or maybe based on whether or not opaque pointers are enabled

There isn't really a meaningful alternative representation for these opaque types when opaque pointers are enabled. So it sounds reasonable to gate it on whether opaque pointers are enabled.

clang/include/clang-c/Index.h

30

I suppose you need to bump CINDEX_VERSION_MINOR for the enum additions?

llvm/docs/SPIRVUsage.rst

103

[aaron.ballman]

clang/include/clang-c/Index.h

30

Correct.

[Anastasia]

Can you link relevant LLVM reviews or documentation and perhaps add some more details into the description?

It's a bit suboptimal though to emit this form of types as SPIR-V specific. Ideally Clang should worry as less as possible about target specifics. So it would be better if we emit OpenCL opaque types as opaque in LLVM IR and then let backends expand those to whatever they need them to be. There isn't technically anything SPIR-V specific in `target("spirv.Image", void, 1, 1, 0, 0, 0, 0, 0)` as this form is just a better representation of the type with more HL information preserved that can either be used or discarded by the optimiser/backend.

My worry is that with the current approach more complexity is being added into the common parts of the compiler. At least it might be better to hide this alternative generation paths into the target implementation itself just like for example we ask target hooks to provide mapping to address spaces. Then we could just have a default path from which targets can derive from...

clang/lib/CodeGen/CGExprScalar.cpp

2263 ↗

(On Diff #486354)

This case at least deserve a comment with explanation. Otherwise I am confused why we end up with a non-pointer type as destination in NullToPointer cast?

clang/lib/CodeGen/CGOpenCLRuntime.cpp

38

Perhaps it's best to split into separate functions and reflect in naming what they correspond to.

58

Why does this need special handling?

[yaxunl]

clang/test/CodeGenOpenCL/cl20-device-side-enqueue.cl

1–6

need to add a non-spir target, otherwise we lose coverage for other targets.

clang/test/CodeGenOpenCL/opencl_types.cl

1–2

need a non-spir target

clang/test/CodeGenOpenCL/sampler.cl

2

need a non-spir target

[jcranmer-intel]

clang/lib/CodeGen/CGOpenCLRuntime.cpp

38

I've decided to pull this out into a new subclass of CGOpenCLRuntime that's target specific.

58

The EXT_OPAQUE_TYPE macro doesn't map cleanly to the SPIR-V OpType* names--the parameters you'd get would be intel_sub_group_avc_mce_payload_t, OCLIntelSubgroupAVCMcePayload, cl_intel_device_side_avc_motion_estimation whereas the SPIR-V type name is supposed to be AvcMcePayloadINTEL. Using the INTEL_SUBGROUP_AVC_TYPE instead allows me to construct the correct SPIR-V type name without having to manually write out each case statement, or adjust macros used in many, many more places. (As it is, I need to rename 4 of the types).

[jcranmer-intel]

Fix some of the code review comments.

[Anastasia]

clang/lib/CodeGen/CGExprScalar.cpp

2260 ↗

(On Diff #487960)

Ok, yet this looks strange to me... do you have an example that hits this code?

At some point we added CK_ZeroToOCLOpaqueType so I wonder if we should be using this instead of CK_NullToPointer here i.e. ideally clang should not assume too early how type are mapped into target specific representation.

clang/lib/CodeGen/CGOpenCLRuntime.cpp

209

Any reason for this? Can we create a constexpr map or enum type that would contain those numbers instead of using hard coded ones scattered around?

235

This looks good in general but we keep CodeGen hierarchy target agnostic so the way to add specifics of an individual target is to extend TargetCodeGenInfo adding a new target hook member function for example it could be getOpenCLSpecificType or something like that. Then you can add SPIR-V specific logic into either CommonSPIRTargetCodeGenInfo or SPIRVTargetCodeGenInfo subclass. I would recommend to look at
createEnqueuedBlockKernel for inspiration https://clang.llvm.org/doxygen/CodeGen_2TargetInfo_8cpp_source.html#l12407.

clang/lib/CodeGen/CodeGenModule.cpp

239 ↗

(On Diff #487960)

Do we want to change old SPIR representation or keep it as is? It seems that SPIR spec defined them as LLVM's opaque pointer types... but I appreciate that for maintenance purposes it's easier to keep those in sync.

llvm/docs/SPIRVUsage.rst

103

Any reference to the document/section explaining these instructions would be useful here.

[jcranmer-intel]

This updates code, and rebases tests on top of trunk.

Note: test issues still haven't been fixed, will fix that likely in the next
revision of the patch.

[jcranmer-intel]

clang/lib/CodeGen/CGExprScalar.cpp

2260 ↗

(On Diff #487960)

I thought I had a test case that caused this code to be executed it, but after removing it and trying out the entire testsuite with it disabled, it never fired. So these changes are indeed unnecessary.

clang/lib/CodeGen/CGOpenCLRuntime.cpp

209

There are constants in the SPIR-V backend that could conceivably be reused (generated from tablegen), but there's no guarantee that we're being compiled with the LLVM SPIR-V backend enabled, which makes it hard to reuse them.

I can add some more comments to explicitly give the references that specify what the values mean, but I'm not entirely certain it's worth building enums just for this one function to use.

clang/lib/CodeGen/CodeGenModule.cpp

239 ↗

(On Diff #487960)

From what I can observe, SPIR and SPIRV seem to be used interchangeably as far as the toolchains are concerned--anything targetting SPIR ends up going through the SPIRV toolchain. This is why I've triggered it for both of the toolchains.

[sidorovd]

clang/lib/CodeGen/TargetInfo.cpp

10978

There probably can be more types

[Anastasia]

clang/lib/CodeGen/CGOpenCLRuntime.cpp

40

I am not sure if this can be called a generic implementation as it has always been a workaround to compensate for absence of representation in LLVM IR... we should probably move such code into default definition of TargetInfo::getOpenCLType...

clang/lib/CodeGen/TargetInfo.cpp

10984

I think the list initialization doesn't do what you are trying to achieve with SmallVector as it appends the elements. You should probably just be using constructor with size i.e. IntParams(6)?

10987

Is this URL likely to stay the same? We normally just add a spec section number e.g.

SPIR-V spec v1.6 rev2 s3.8.

10988

We should avoid adding file paths and names as they can change without updating the comment.

10990

Ok, is the order of parameters defined or documented somewhere? Would it make sense to create some kind of a local enum map containing the indices and use enum members instead of constants to improve readability/maintenance?

11013

Ok, so the old SPIR format will change too!

I don't have any objections but want to make sure it's not accidental.

[jcranmer-intel]

Mostly test fixes.

clang/test/CodeGenOpenCL/opencl_types.cl

1–2

This file checks AMD already.

[iliya-diyachkov]

clang/lib/CodeGen/CGOpenCLRuntime.cpp

100–113

Perhaps use early exits like this, or even removing 'else' clause.

[jcranmer-intel]

Friendly review ping?

[jcranmer-intel]

Documentation tweaks.

[jcranmer-intel]

clang/lib/CodeGen/TargetInfo.cpp

10984

The list initialization is filling in 6 0 elements, which is intentional.

10990

They're documented in the link given in the first comment in this method: https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpTypeImage

[jcranmer-intel]

Friendly review ping?

[Anastasia]

LGTM! Thanks

clang/lib/CodeGen/TargetInfo.h

383

The default behaviour doesn't return a target extension type so maybe it's better to change the comment to:

Return an LLVM type that corresponds to an OpenCL type

clang/include/clang-c/Index.h

 #include "clang-c/CXFile.h"
 #include "clang-c/CXSourceLocation.h"
 #include "clang-c/CXString.h"
 #include "clang-c/ExternC.h"
 #include "clang-c/Platform.h"
 
 /**
  * The version constants for the libclang API.
  * CINDEX_VERSION_MINOR should increase when there are API additions.
  * CINDEX_VERSION_MAJOR is intended for "major" source/ABI breaking changes.
  *
  * The policy about the libclang API was always to keep it source and ABI
  * compatible, thus CINDEX_VERSION_MAJOR is expected to remain stable.
  */
 #define CINDEX_VERSION_MAJOR 0
 #define CINDEX_VERSION_MINOR 64
 
enum CXTypeKind {
   CXType_OCLIntelSubgroupAVCMcePayload = 164,
   CXType_OCLIntelSubgroupAVCImePayload = 165,
   CXType_OCLIntelSubgroupAVCRefPayload = 166,
   CXType_OCLIntelSubgroupAVCSicPayload = 167,
   CXType_OCLIntelSubgroupAVCMceResult = 168,
   CXType_OCLIntelSubgroupAVCImeResult = 169,
   CXType_OCLIntelSubgroupAVCRefResult = 170,
   CXType_OCLIntelSubgroupAVCSicResult = 171,
+  CXType_OCLIntelSubgroupAVCImeResultSingleReferenceStreamout = 172,
+  CXType_OCLIntelSubgroupAVCImeResultDualReferenceStreamout = 173,
+  CXType_OCLIntelSubgroupAVCImeSingleReferenceStreamin = 174,
+  CXType_OCLIntelSubgroupAVCImeDualReferenceStreamin = 175,
+
+  /* Old aliases for AVC OpenCL extension types. */
   CXType_OCLIntelSubgroupAVCImeResultSingleRefStreamout = 172,
   CXType_OCLIntelSubgroupAVCImeResultDualRefStreamout = 173,
   CXType_OCLIntelSubgroupAVCImeSingleRefStreamin = 174,
-
   CXType_OCLIntelSubgroupAVCImeDualRefStreamin = 175,
 
   CXType_ExtVector = 176,
   CXType_Atomic = 177,
   CXType_BTFTagAttributed = 178
 };
 
 /**

clang/include/clang/Basic/OpenCLExtensionTypes.def

 INTEL_SUBGROUP_AVC_TYPE(mce_payload_t, McePayload)
 INTEL_SUBGROUP_AVC_TYPE(ime_payload_t, ImePayload)
 INTEL_SUBGROUP_AVC_TYPE(ref_payload_t, RefPayload)
 INTEL_SUBGROUP_AVC_TYPE(sic_payload_t, SicPayload)
 INTEL_SUBGROUP_AVC_TYPE(mce_result_t, MceResult)
 INTEL_SUBGROUP_AVC_TYPE(ime_result_t, ImeResult)
 INTEL_SUBGROUP_AVC_TYPE(ref_result_t, RefResult)
 INTEL_SUBGROUP_AVC_TYPE(sic_result_t, SicResult)
-INTEL_SUBGROUP_AVC_TYPE(ime_result_single_reference_streamout_t, ImeResultSingleRefStreamout)
-INTEL_SUBGROUP_AVC_TYPE(ime_result_dual_reference_streamout_t, ImeResultDualRefStreamout)
-INTEL_SUBGROUP_AVC_TYPE(ime_single_reference_streamin_t, ImeSingleRefStreamin)
-INTEL_SUBGROUP_AVC_TYPE(ime_dual_reference_streamin_t, ImeDualRefStreamin)
+INTEL_SUBGROUP_AVC_TYPE(ime_result_single_reference_streamout_t, ImeResultSingleReferenceStreamout)
+INTEL_SUBGROUP_AVC_TYPE(ime_result_dual_reference_streamout_t, ImeResultDualReferenceStreamout)
+INTEL_SUBGROUP_AVC_TYPE(ime_single_reference_streamin_t, ImeSingleReferenceStreamin)
+INTEL_SUBGROUP_AVC_TYPE(ime_dual_reference_streamin_t, ImeDualReferenceStreamin)
 
 #undef INTEL_SUBGROUP_AVC_TYPE
 #endif // INTEL_SUBGROUP_AVC_TYPE
 
 #undef EXT_OPAQUE_TYPE
 

clang/lib/CodeGen/CGOpenCLRuntime.h

 class CodeGenFunction;
 class CodeGenModule;
 
 class CGOpenCLRuntime {
 protected:
   CodeGenModule &CGM;
   llvm::Type *PipeROTy;
   llvm::Type *PipeWOTy;
-  llvm::PointerType *SamplerTy;
+  llvm::Type *SamplerTy;
   llvm::StringMap<llvm::PointerType *> CachedTys;
 
   /// Structure for enqueued block information.
   struct EnqueuedBlockInfo {
     llvm::Function *InvokeFunc; /// Block invoke function.
     llvm::Value *KernelHandle;  /// Enqueued block kernel reference.
     llvm::Value *BlockArg;      /// The first argument to enqueued block kernel.
     llvm::Type *BlockTy;        /// Type of the block argument.
public:
   /// CodeGenFunction::EmitStaticVarDecl to emit an internal global for D.
   virtual void EmitWorkGroupLocalVarDecl(CodeGenFunction &CGF,
                                          const VarDecl &D);
 
   virtual llvm::Type *convertOpenCLSpecificType(const Type *T);
 
   virtual llvm::Type *getPipeType(const PipeType *T);
 
-  llvm::PointerType *getSamplerType(const Type *T);
+  llvm::Type *getSamplerType(const Type *T);
 
   // Returns a value which indicates the size in bytes of the pipe
   // element.
   virtual llvm::Value *getPipeElemSize(const Expr *PipeArg);
 
   // Returns a value which indicates the alignment in bytes of the pipe
   // element.
   virtual llvm::Value *getPipeElemAlign(const Expr *PipeArg);

clang/lib/CodeGen/CGOpenCLRuntime.cpp

 CGOpenCLRuntime::~CGOpenCLRuntime() {}
 
 void CGOpenCLRuntime::EmitWorkGroupLocalVarDecl(CodeGenFunction &CGF,
                                                 const VarDecl &D) {
   return CGF.EmitStaticVarDecl(D, llvm::GlobalValue::InternalLinkage);
 }
 
 llvm::Type *CGOpenCLRuntime::convertOpenCLSpecificType(const Type *T) {
-  assert(T->isOpenCLSpecificType() &&
-         "Not an OpenCL specific type!");
+  assert(T->isOpenCLSpecificType() && "Not an OpenCL specific type!");
+
+  // Check if the target has a specific translation for this type first.
+  if (llvm::Type *TransTy = CGM.getTargetCodeGenInfo().getOpenCLType(CGM, T))
+    return TransTy;
 
   switch (cast<BuiltinType>(T)->getKind()) {
   default:
     llvm_unreachable("Unexpected opencl builtin type!");
     return nullptr;
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
   case BuiltinType::Id:                                                        \
     return getPointerType(T, "opencl." #ImgType "_" #Suffix "_t");
 #include "clang/Basic/OpenCLImageTypes.def"
   case BuiltinType::OCLSampler:
     return getSamplerType(T);
   case BuiltinType::OCLEvent:
     return getPointerType(T, "opencl.event_t");
   case BuiltinType::OCLClkEvent:
     return getPointerType(T, "opencl.clk_event_t");
   case BuiltinType::OCLQueue:
     return getPointerType(T, "opencl.queue_t");
   case BuiltinType::OCLReserveID:
     return getPointerType(T, "opencl.reserve_id_t");
 #define EXT_OPAQUE_TYPE(ExtType, Id, Ext)                                      \
   case BuiltinType::Id:                                                        \
     return getPointerType(T, "opencl." #ExtType);
 #include "clang/Basic/OpenCLExtensionTypes.def"
   }
 }
 
 llvm::PointerType *CGOpenCLRuntime::getPointerType(const Type *T,
                                                    StringRef Name) {
   auto I = CachedTys.find(Name);
   if (I != CachedTys.end())
     return I->second;
 
   llvm::LLVMContext &Ctx = CGM.getLLVMContext();
   uint32_t AddrSpc = CGM.getContext().getTargetAddressSpace(
       CGM.getContext().getOpenCLTypeAddrSpace(T));
   auto *PTy =
       llvm::PointerType::get(llvm::StructType::create(Ctx, Name), AddrSpc);
   CachedTys[Name] = PTy;
   return PTy;
 }
 
 llvm::Type *CGOpenCLRuntime::getPipeType(const PipeType *T) {
+  if (llvm::Type *PipeTy = CGM.getTargetCodeGenInfo().getOpenCLType(CGM, T))
+    return PipeTy;
+
   if (T->isReadOnly())
     return getPipeType(T, "opencl.pipe_ro_t", PipeROTy);
   else
     return getPipeType(T, "opencl.pipe_wo_t", PipeWOTy);
 }
 
 llvm::Type *CGOpenCLRuntime::getPipeType(const PipeType *T, StringRef Name,
                                          llvm::Type *&PipeTy) {
   if (!PipeTy)
     PipeTy = llvm::PointerType::get(llvm::StructType::create(
       CGM.getLLVMContext(), Name),
       CGM.getContext().getTargetAddressSpace(
           CGM.getContext().getOpenCLTypeAddrSpace(T)));
   return PipeTy;
 }
 
-llvm::PointerType *CGOpenCLRuntime::getSamplerType(const Type *T) {
-  if (!SamplerTy)
-    SamplerTy = llvm::PointerType::get(llvm::StructType::create(
-      CGM.getLLVMContext(), "opencl.sampler_t"),
+llvm::Type *CGOpenCLRuntime::getSamplerType(const Type *T) {
+  if (SamplerTy)
+    return SamplerTy;
+
+  if (llvm::Type *TransTy = CGM.getTargetCodeGenInfo().getOpenCLType(
+          CGM, CGM.getContext().OCLSamplerTy.getTypePtr()))
+    SamplerTy = TransTy;
+  else
+    SamplerTy = llvm::PointerType::get(
+        llvm::StructType::create(CGM.getLLVMContext(), "opencl.sampler_t"),
         CGM.getContext().getTargetAddressSpace(
             CGM.getContext().getOpenCLTypeAddrSpace(T)));
   return SamplerTy;
 }
 
 llvm::Value *CGOpenCLRuntime::getPipeElemSize(const Expr *PipeArg) {
   const PipeType *PipeTy = PipeArg->getType()->castAs<PipeType>();
   // The type of the last (implicit) argument to be passed.
   llvm::Type *Int32Ty = llvm::IntegerType::getInt32Ty(CGM.getLLVMContext());
   unsigned TypeSize = CGM.getContext()
                           .getTypeSizeInChars(PipeTy->getElementType())
                           .getQuantity();
CGOpenCLRuntime::emitOpenCLEnqueuedBlock(CodeGenFunction &CGF, const Expr *E) {
   }
 
   auto *F = CGF.getTargetHooks().createEnqueuedBlockKernel(
       CGF, EnqueuedBlockMap[Block].InvokeFunc, EnqueuedBlockMap[Block].BlockTy);
 
   // The common part of the post-processing of the kernel goes here.
   EnqueuedBlockMap[Block].KernelHandle = F;
   return EnqueuedBlockMap[Block];
 }

clang/lib/CodeGen/TargetInfo.h

public:
   }
   /// Emit the device-side copy of the builtin texture type.
   virtual bool emitCUDADeviceBuiltinTextureDeviceCopy(CodeGenFunction &CGF,
                                                       LValue Dst,
                                                       LValue Src) const {
     // DO NOTHING by default.
     return false;
   }
+
+  /// Return an LLVM type that corresponds to an OpenCL type.
+  virtual llvm::Type *getOpenCLType(CodeGenModule &CGM, const Type *T) const {
+    return nullptr;
+  }
 };
 
 } // namespace CodeGen
 } // namespace clang
 
 #endif // LLVM_CLANG_LIB_CODEGEN_TARGETINFO_H

clang/lib/CodeGen/TargetInfo.cpp

CommonSPIRTargetCodeGenInfo(std::unique_ptr<ABIInfo> ABIInfo)
       : TargetCodeGenInfo(std::move(ABIInfo)) {}
 
   LangAS getASTAllocaAddressSpace() const override {
     return getLangASFromTargetAS(
         getABIInfo().getDataLayout().getAllocaAddrSpace());
   }
 
   unsigned getOpenCLKernelCallingConv() const override;
+  llvm::Type *getOpenCLType(CodeGenModule &CGM, const Type *T) const override;
 };
 class SPIRVTargetCodeGenInfo : public CommonSPIRTargetCodeGenInfo {
 public:
   SPIRVTargetCodeGenInfo(CodeGen::CodeGenTypes &CGT)
       : CommonSPIRTargetCodeGenInfo(std::make_unique<SPIRVABIInfo>(CGT)) {}
   void setCUDAKernelCallingConvention(const FunctionType *&FT) const override;
 };
 } // End anonymous namespace.
if (const ArrayType *AT = QT->getAsArrayTypeUnsafe()) {
       // Thus we don't call appendQualifier() here.
       return appendArrayType(Enc, QT, AT, CGM, TSC, "*");
     }
     return appendType(Enc, QT, CGM, TSC);
   }
   return false;
 }
 
+/// Construct a SPIR-V target extension type for the given OpenCL image type.
+static llvm::Type *getSPIRVImageType(llvm::LLVMContext &Ctx, StringRef BaseType,
+                                     StringRef OpenCLName,
+                                     unsigned AccessQualifier) {
+  // These parameters compare to the operands of OpTypeImage (see
+  // https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpTypeImage
+  // for more details). The first 6 integer parameters all default to 0, and
+  // will be changed to 1 only for the image type(s) that set the parameter to
+  // one. The 7th integer parameter is the access qualifier, which is tacked on
+  // at the end.
+  SmallVector<unsigned, 7> IntParams = {0, 0, 0, 0, 0, 0};
+
+  // Choose the dimension of the image--this corresponds to the Dim enum in
+  // SPIR-V (first integer parameter of OpTypeImage).
+  if (OpenCLName.startswith("image2d"))
+    IntParams[0] = 1; // 1D
+  else if (OpenCLName.startswith("image3d"))
+    IntParams[0] = 2; // 2D
+  else if (OpenCLName == "image1d_buffer")
+    IntParams[0] = 5; // Buffer
+  else
+    assert(OpenCLName.startswith("image1d") && "Unknown image type");
+
+  // Set the other integer parameters of OpTypeImage if necessary. Note that the
+  // OpenCL image types don't provide any information for the Sampled or
+  // Image Format parameters.
+  if (OpenCLName.contains("_depth"))
+    IntParams[1] = 1;
+  if (OpenCLName.contains("_array"))
+    IntParams[2] = 1;
+  if (OpenCLName.contains("_msaa"))
+    IntParams[3] = 1;
+
+  // Access qualifier
+  IntParams.push_back(AccessQualifier);
+
+  return llvm::TargetExtType::get(Ctx, BaseType, {llvm::Type::getVoidTy(Ctx)},
+                                  IntParams);
+}
+
+llvm::Type *CommonSPIRTargetCodeGenInfo::getOpenCLType(CodeGenModule &CGM,
+                                                       const Type *Ty) const {
+  llvm::LLVMContext &Ctx = CGM.getLLVMContext();
+  if (auto *PipeTy = dyn_cast<PipeType>(Ty))
+    return llvm::TargetExtType::get(Ctx, "spirv.Pipe", {},
+                                    {!PipeTy->isReadOnly()});
+  if (auto *BuiltinTy = dyn_cast<BuiltinType>(Ty)) {
+    enum AccessQualifier : unsigned { AQ_ro = 0, AQ_wo = 1, AQ_rw = 2 };
+    switch (BuiltinTy->getKind()) {
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+    case BuiltinType::Id:                                                      \
+      return getSPIRVImageType(Ctx, "spirv.Image", #ImgType, AQ_##Suffix);
+#include "clang/Basic/OpenCLImageTypes.def"
+    case BuiltinType::OCLSampler:
+      return llvm::TargetExtType::get(Ctx, "spirv.Sampler");
+    case BuiltinType::OCLEvent:
+      return llvm::TargetExtType::get(Ctx, "spirv.Event");
+    case BuiltinType::OCLClkEvent:
+      return llvm::TargetExtType::get(Ctx, "spirv.DeviceEvent");
+    case BuiltinType::OCLQueue:
+      return llvm::TargetExtType::get(Ctx, "spirv.Queue");
+    case BuiltinType::OCLReserveID:
+      return llvm::TargetExtType::get(Ctx, "spirv.ReserveId");
+#define INTEL_SUBGROUP_AVC_TYPE(Name, Id)                                      \
+    case BuiltinType::OCLIntelSubgroupAVC##Id:                                 \
+      return llvm::TargetExtType::get(Ctx, "spirv.Avc" #Id "INTEL");
+#include "clang/Basic/OpenCLExtensionTypes.def"
+    default:
+      return nullptr;
+    }
+  }
+
+  return nullptr;
+}
 //===----------------------------------------------------------------------===//
 // RISCV ABI Implementation
 //===----------------------------------------------------------------------===//
 
 namespace {
 class RISCVABIInfo : public DefaultABIInfo {
 private:
   // Size of the integer ('x') registers in bits.

clang/test/CodeGenOpenCL/cast_image.cl

 // RUN: %clang_cc1 -no-opaque-pointers -emit-llvm -o - -triple amdgcn--amdhsa %s | FileCheck --check-prefix=AMDGCN %s
-// RUN: %clang_cc1 -no-opaque-pointers -emit-llvm -o - -triple spir-unknown-unknown %s | FileCheck --check-prefix=SPIR %s
+// RUN: %clang_cc1 -no-opaque-pointers -emit-llvm -o - -triple x86_64-unknown-unknown %s | FileCheck --check-prefix=X86 %s
 
 #ifdef __AMDGCN__
 
 constant int* convert(image2d_t img) {
   // AMDGCN: bitcast %opencl.image2d_ro_t addrspace(4)* %img to i32 addrspace(4)*
   return __builtin_astype(img, constant int*);
 }
 
 #else
 
 global int* convert(image2d_t img) {
-  // SPIR: bitcast %opencl.image2d_ro_t addrspace(1)* %img to i32 addrspace(1)*
+  // X86: bitcast %opencl.image2d_ro_t* %img to i32*
   return __builtin_astype(img, global int*);
 }
 
 #endif

clang/test/CodeGenOpenCL/cl20-device-side-enqueue-attributes.cl

 }
 // SPIR32: Function Attrs: convergent noinline norecurse nounwind optnone
 // SPIR32-LABEL: define {{[^@]+}}@device_side_enqueue
 // SPIR32-SAME: (ptr addrspace(1) align 4 [[A:%.*]], ptr addrspace(1) align 4 [[B:%.*]], i32 [[I:%.*]]) #[[ATTR0:[0-9]+]] !kernel_arg_addr_space !2 !kernel_arg_access_qual !3 !kernel_arg_type !4 !kernel_arg_base_type !4 !kernel_arg_type_qual !5 {
 // SPIR32-NEXT:  entry:
 // SPIR32-NEXT:    [[A_ADDR:%.*]] = alloca ptr addrspace(1), align 4
 // SPIR32-NEXT:    [[B_ADDR:%.*]] = alloca ptr addrspace(1), align 4
 // SPIR32-NEXT:    [[I_ADDR:%.*]] = alloca i32, align 4
-// SPIR32-NEXT:    [[DEFAULT_QUEUE:%.*]] = alloca ptr, align 4
+// SPIR32-NEXT:    [[DEFAULT_QUEUE:%.*]] = alloca target("spirv.Queue"), align 4
 // SPIR32-NEXT:    [[FLAGS:%.*]] = alloca i32, align 4
 // SPIR32-NEXT:    [[NDRANGE:%.*]] = alloca [[STRUCT_NDRANGE_T:%.*]], align 4
 // SPIR32-NEXT:    [[TMP:%.*]] = alloca [[STRUCT_NDRANGE_T]], align 4
 // SPIR32-NEXT:    [[BLOCK:%.*]] = alloca <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, align 4
 // SPIR32-NEXT:    store ptr addrspace(1) [[A]], ptr [[A_ADDR]], align 4
 // SPIR32-NEXT:    store ptr addrspace(1) [[B]], ptr [[B_ADDR]], align 4
 // SPIR32-NEXT:    store i32 [[I]], ptr [[I_ADDR]], align 4
 // SPIR32-NEXT:    store i32 0, ptr [[FLAGS]], align 4
-// SPIR32-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DEFAULT_QUEUE]], align 4
+// SPIR32-NEXT:    [[TMP0:%.*]] = load target("spirv.Queue"), ptr [[DEFAULT_QUEUE]], align 4
 // SPIR32-NEXT:    [[TMP1:%.*]] = load i32, ptr [[FLAGS]], align 4
 // SPIR32-NEXT:    call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[TMP]], ptr align 4 [[NDRANGE]], i32 4, i1 false)
 // SPIR32-NEXT:    [[BLOCK_SIZE:%.*]] = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr [[BLOCK]], i32 0, i32 0
 // SPIR32-NEXT:    store i32 24, ptr [[BLOCK_SIZE]], align 4
 // SPIR32-NEXT:    [[BLOCK_ALIGN:%.*]] = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr [[BLOCK]], i32 0, i32 1
 // SPIR32-NEXT:    store i32 4, ptr [[BLOCK_ALIGN]], align 4
 // SPIR32-NEXT:    [[BLOCK_INVOKE:%.*]] = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr [[BLOCK]], i32 0, i32 2
 // SPIR32-NEXT:    store ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke to ptr addrspace(4)), ptr [[BLOCK_INVOKE]], align 4
 // SPIR32-NEXT:    [[BLOCK_CAPTURED:%.*]] = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr [[BLOCK]], i32 0, i32 3
 // SPIR32-NEXT:    [[TMP2:%.*]] = load ptr addrspace(1), ptr [[A_ADDR]], align 4
 // SPIR32-NEXT:    store ptr addrspace(1) [[TMP2]], ptr [[BLOCK_CAPTURED]], align 4
 // SPIR32-NEXT:    [[BLOCK_CAPTURED1:%.*]] = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr [[BLOCK]], i32 0, i32 4
 // SPIR32-NEXT:    [[TMP3:%.*]] = load i32, ptr [[I_ADDR]], align 4
 // SPIR32-NEXT:    store i32 [[TMP3]], ptr [[BLOCK_CAPTURED1]], align 4
 // SPIR32-NEXT:    [[BLOCK_CAPTURED2:%.*]] = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr [[BLOCK]], i32 0, i32 5
 // SPIR32-NEXT:    [[TMP4:%.*]] = load ptr addrspace(1), ptr [[B_ADDR]], align 4
 // SPIR32-NEXT:    store ptr addrspace(1) [[TMP4]], ptr [[BLOCK_CAPTURED2]], align 4
 // SPIR32-NEXT:    [[TMP5:%.*]] = addrspacecast ptr [[BLOCK]] to ptr addrspace(4)
-// SPIR32-NEXT:    [[TMP6:%.*]] = call spir_func i32 @__enqueue_kernel_basic(ptr [[TMP0]], i32 [[TMP1]], ptr byval([[STRUCT_NDRANGE_T]]) [[TMP]], ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_kernel to ptr addrspace(4)), ptr addrspace(4) [[TMP5]])
+// SPIR32-NEXT:    [[TMP6:%.*]] = call spir_func i32 @__enqueue_kernel_basic(target("spirv.Queue") [[TMP0]], i32 [[TMP1]], ptr byval([[STRUCT_NDRANGE_T]]) [[TMP]], ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_kernel to ptr addrspace(4)), ptr addrspace(4) [[TMP5]])
 // SPIR32-NEXT:    ret void
 //
 //
 // SPIR32: Function Attrs: convergent noinline nounwind optnone
 // SPIR32-LABEL: define {{[^@]+}}@__device_side_enqueue_block_invoke
 // SPIR32-SAME: (ptr addrspace(4) [[DOTBLOCK_DESCRIPTOR:%.*]]) #[[ATTR2:[0-9]+]] {
 // SPIR32-NEXT:  entry:
 // SPIR32-NEXT:    [[DOTBLOCK_DESCRIPTOR_ADDR:%.*]] = alloca ptr addrspace(4), align 4

 //
 // STRICTFP: Function Attrs: convergent noinline norecurse nounwind optnone strictfp
 // STRICTFP-LABEL: define {{[^@]+}}@device_side_enqueue
 // STRICTFP-SAME: (ptr addrspace(1) align 4 [[A:%.*]], ptr addrspace(1) align 4 [[B:%.*]], i32 [[I:%.*]]) #[[ATTR0:[0-9]+]] !kernel_arg_addr_space !2 !kernel_arg_access_qual !3 !kernel_arg_type !4 !kernel_arg_base_type !4 !kernel_arg_type_qual !5 {
 // STRICTFP-NEXT:  entry:
 // STRICTFP-NEXT:    [[A_ADDR:%.*]] = alloca ptr addrspace(1), align 4
 // STRICTFP-NEXT:    [[B_ADDR:%.*]] = alloca ptr addrspace(1), align 4
 // STRICTFP-NEXT:    [[I_ADDR:%.*]] = alloca i32, align 4
-// STRICTFP-NEXT:    [[DEFAULT_QUEUE:%.*]] = alloca ptr, align 4
+// STRICTFP-NEXT:    [[DEFAULT_QUEUE:%.*]] = alloca target("spirv.Queue"), align 4
 // STRICTFP-NEXT:    [[FLAGS:%.*]] = alloca i32, align 4
 // STRICTFP-NEXT:    [[NDRANGE:%.*]] = alloca [[STRUCT_NDRANGE_T:%.*]], align 4
 // STRICTFP-NEXT:    [[TMP:%.*]] = alloca [[STRUCT_NDRANGE_T]], align 4
 // STRICTFP-NEXT:    [[BLOCK:%.*]] = alloca <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, align 4
 // STRICTFP-NEXT:    store ptr addrspace(1) [[A]], ptr [[A_ADDR]], align 4
 // STRICTFP-NEXT:    store ptr addrspace(1) [[B]], ptr [[B_ADDR]], align 4
 // STRICTFP-NEXT:    store i32 [[I]], ptr [[I_ADDR]], align 4
 // STRICTFP-NEXT:    store i32 0, ptr [[FLAGS]], align 4
-// STRICTFP-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DEFAULT_QUEUE]], align 4
+// STRICTFP-NEXT:    [[TMP0:%.*]] = load target("spirv.Queue"), ptr [[DEFAULT_QUEUE]], align 4
 // STRICTFP-NEXT:    [[TMP1:%.*]] = load i32, ptr [[FLAGS]], align 4
 // STRICTFP-NEXT:    call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[TMP]], ptr align 4 [[NDRANGE]], i32 4, i1 false) #[[ATTR5:[0-9]+]]
 // STRICTFP-NEXT:    [[BLOCK_SIZE:%.*]] = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr [[BLOCK]], i32 0, i32 0
 // STRICTFP-NEXT:    store i32 24, ptr [[BLOCK_SIZE]], align 4
 // STRICTFP-NEXT:    [[BLOCK_ALIGN:%.*]] = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr [[BLOCK]], i32 0, i32 1
 // STRICTFP-NEXT:    store i32 4, ptr [[BLOCK_ALIGN]], align 4
 // STRICTFP-NEXT:    [[BLOCK_INVOKE:%.*]] = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr [[BLOCK]], i32 0, i32 2
 // STRICTFP-NEXT:    store ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke to ptr addrspace(4)), ptr [[BLOCK_INVOKE]], align 4
 // STRICTFP-NEXT:    [[BLOCK_CAPTURED:%.*]] = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr [[BLOCK]], i32 0, i32 3
 // STRICTFP-NEXT:    [[TMP2:%.*]] = load ptr addrspace(1), ptr [[A_ADDR]], align 4
 // STRICTFP-NEXT:    store ptr addrspace(1) [[TMP2]], ptr [[BLOCK_CAPTURED]], align 4
 // STRICTFP-NEXT:    [[BLOCK_CAPTURED1:%.*]] = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr [[BLOCK]], i32 0, i32 4
 // STRICTFP-NEXT:    [[TMP3:%.*]] = load i32, ptr [[I_ADDR]], align 4
 // STRICTFP-NEXT:    store i32 [[TMP3]], ptr [[BLOCK_CAPTURED1]], align 4
 // STRICTFP-NEXT:    [[BLOCK_CAPTURED2:%.*]] = getelementptr inbounds <{ i32, i32, ptr addrspace(4), ptr addrspace(1), i32, ptr addrspace(1) }>, ptr [[BLOCK]], i32 0, i32 5
 // STRICTFP-NEXT:    [[TMP4:%.*]] = load ptr addrspace(1), ptr [[B_ADDR]], align 4
 // STRICTFP-NEXT:    store ptr addrspace(1) [[TMP4]], ptr [[BLOCK_CAPTURED2]], align 4
 // STRICTFP-NEXT:    [[TMP5:%.*]] = addrspacecast ptr [[BLOCK]] to ptr addrspace(4)
-// STRICTFP-NEXT:    [[TMP6:%.*]] = call spir_func i32 @__enqueue_kernel_basic(ptr [[TMP0]], i32 [[TMP1]], ptr byval([[STRUCT_NDRANGE_T]]) [[TMP]], ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_kernel to ptr addrspace(4)), ptr addrspace(4) [[TMP5]])
+// STRICTFP-NEXT:    [[TMP6:%.*]] = call spir_func i32 @__enqueue_kernel_basic(target("spirv.Queue") [[TMP0]], i32 [[TMP1]], ptr byval([[STRUCT_NDRANGE_T]]) [[TMP]], ptr addrspace(4) addrspacecast (ptr @__device_side_enqueue_block_invoke_kernel to ptr addrspace(4)), ptr addrspace(4) [[TMP5]])
 // STRICTFP-NEXT:    ret void
 //
 //
 // STRICTFP: Function Attrs: convergent noinline nounwind optnone strictfp
 // STRICTFP-LABEL: define {{[^@]+}}@__device_side_enqueue_block_invoke
 // STRICTFP-SAME: (ptr addrspace(4) [[DOTBLOCK_DESCRIPTOR:%.*]]) #[[ATTR2:[0-9]+]] {
 // STRICTFP-NEXT:  entry:
 // STRICTFP-NEXT:    [[DOTBLOCK_DESCRIPTOR_ADDR:%.*]] = alloca ptr addrspace(4), align 4

clang/test/CodeGenOpenCL/cl20-device-side-enqueue.cl

-// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL2.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir-unknown-unknown" | FileCheck %s --check-prefix=COMMON --check-prefix=B32
-// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL2.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir64-unknown-unknown" | FileCheck %s --check-prefix=COMMON --check-prefix=B64
+// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL2.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir-unknown-unknown" | FileCheck %s --check-prefixes=COMMON,B32,SPIR
+// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL2.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir64-unknown-unknown" | FileCheck %s --check-prefixes=COMMON,B64,SPIR
 // RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL2.0 -ffake-address-space-map -O1 -emit-llvm -o - -triple "spir64-unknown-unknown" | FileCheck %s --check-prefix=CHECK-LIFETIMES
-// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL3.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir-unknown-unknown" | FileCheck %s --check-prefix=COMMON --check-prefix=B32
-// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL3.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir64-unknown-unknown" | FileCheck %s --check-prefix=COMMON --check-prefix=B64
+// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL3.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir-unknown-unknown" | FileCheck %s --check-prefixes=COMMON,B32,SPIR
+// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL3.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "spir64-unknown-unknown" | FileCheck %s --check-prefixes=COMMON,B64,SPIR
 // RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL3.0 -ffake-address-space-map -O1 -emit-llvm -o - -triple "spir64-unknown-unknown" | FileCheck %s --check-prefix=CHECK-LIFETIMES
+// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL2.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "x86_64-unknown-linux-gnu" | FileCheck %s --check-prefixes=COMMON,B64,X86
+// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL3.0 -ffake-address-space-map -O0 -emit-llvm -o - -triple "x86_64-unknown-linux-gnu" | FileCheck %s --check-prefixes=COMMON,B64,X86
+// RUN: %clang_cc1 -no-enable-noundef-analysis %s -cl-std=CL3.0 -ffake-address-space-map -O1 -emit-llvm -o - -triple "x86_64-unknown-linux-gnu" | FileCheck %s --check-prefix=CHECK-LIFETIMES
 
 #pragma OPENCL EXTENSION cl_khr_subgroups : enable
 
 typedef void (^bl_t)(local void *);
 typedef struct {int a;} ndrange_t;
 
 // For a block global variable, first emit the block literal as a global variable, then emit the block variable itself.
 // COMMON: [[BL_GLOBAL:@__block_literal_global[^ ]*]] = internal addrspace(1) constant { i32, i32, ptr addrspace(4) } { i32 {{[0-9]+}}, i32 {{[0-9]+}}, ptr addrspace(4) addrspacecast (ptr [[INV_G:@[^ ]+]] to ptr addrspace(4)) }

 // COMMON: [[BLG6:@__block_literal_global[^ ]*]] = internal addrspace(1) constant { i32, i32, ptr addrspace(4) } { i32 {{[0-9]+}}, i32 {{[0-9]+}}, ptr addrspace(4) addrspacecast (ptr {{@[^ ]+}} to ptr addrspace(4)) }
 // COMMON: [[BLG7:@__block_literal_global[^ ]*]] = internal addrspace(1) constant { i32, i32, ptr addrspace(4) } { i32 {{[0-9]+}}, i32 {{[0-9]+}}, ptr addrspace(4) addrspacecast (ptr {{@[^ ]+}} to ptr addrspace(4)) }
 // COMMON: [[BLG8:@__block_literal_global[^ ]*]] = internal addrspace(1) constant { i32, i32, ptr addrspace(4) } { i32 {{[0-9]+}}, i32 {{[0-9]+}}, ptr addrspace(4) addrspacecast (ptr [[INVG8:@[^ ]+]] to ptr addrspace(4)) }
 // COMMON: [[BLG9:@__block_literal_global[^ ]*]] = internal addrspace(1) constant { i32, i32, ptr addrspace(4) } { i32 {{[0-9]+}}, i32 {{[0-9]+}}, ptr addrspace(4) addrspacecast (ptr [[INVG9:@[^ ]+]] to ptr addrspace(4)) }
 // COMMON: [[BLG10:@__block_literal_global[^ ]*]] = internal addrspace(1) constant { i32, i32, ptr addrspace(4) } { i32 {{[0-9]+}}, i32 {{[0-9]+}}, ptr addrspace(4) addrspacecast (ptr {{@[^ ]+}} to ptr addrspace(4)) }
 // COMMON: [[BLG11:@__block_literal_global[^ ]*]] = internal addrspace(1) constant { i32, i32, ptr addrspace(4) } { i32 {{[0-9]+}}, i32 {{[0-9]+}}, ptr addrspace(4) addrspacecast (ptr {{@[^ ]+}} to ptr addrspace(4)) }
 
 // Emits block literal [[BL_GLOBAL]], invoke function [[INV_G]] and global block variable @block_G
-// COMMON: define internal spir_func void [[INV_G]](ptr addrspace(4) %{{.*}}, ptr addrspace(3) %{{.*}})
+// COMMON: define internal {{(spir_func )?}}void [[INV_G]](ptr addrspace(4) %{{.*}}, ptr addrspace(3) %{{.*}})
 const bl_t block_G = (bl_t) ^ (local void *a) {};
 
 void callee(int id, __global int *out) {
   out[id] = id;
 }
 
 // COMMON-LABEL: define{{.*}} spir_kernel void @device_side_enqueue(ptr addrspace(1) align 4 %{{.*}}, ptr addrspace(1) align 4 %b, i32 %i)
 kernel void device_side_enqueue(global int *a, global int *b, int i) {
-  // COMMON: %default_queue = alloca ptr
+  // SPIR: %default_queue = alloca target("spirv.Queue")
+  // X86: %default_queue = alloca ptr
   queue_t default_queue;
   // COMMON: %flags = alloca i32
   unsigned flags = 0;
   // COMMON: %ndrange = alloca %struct.ndrange_t
   ndrange_t ndrange;
-  // COMMON: %clk_event = alloca ptr
+  // SPIR: %clk_event = alloca target("spirv.DeviceEvent")
+  // X86: %clk_event = alloca ptr
   clk_event_t clk_event;
-  // COMMON: %event_wait_list = alloca ptr
+  // SPIR: %event_wait_list = alloca target("spirv.DeviceEvent")
+  // X86: %event_wait_list = alloca ptr
   clk_event_t event_wait_list;
-  // COMMON: %event_wait_list2 = alloca [1 x ptr]
+  // SPIR: %event_wait_list2 = alloca [1 x target("spirv.DeviceEvent")]
+  // X86: %event_wait_list2 = alloca [1 x ptr]
   clk_event_t event_wait_list2[] = {clk_event};
 
   // COMMON: [[NDR:%[a-z0-9]+]] = alloca %struct.ndrange_t, align 4
 
   // B32: %[[BLOCK_SIZES1:.*]] = alloca [1 x i32]
   // B64: %[[BLOCK_SIZES1:.*]] = alloca [1 x i64]
   // CHECK-LIFETIMES: %[[BLOCK_SIZES1:.*]] = alloca [1 x i64]
   // B32: %[[BLOCK_SIZES2:.*]] = alloca [1 x i32]
kernel void device_side_enqueue(global int *a, global int *b, int i) {
   // B32: %[[BLOCK_SIZES6:.*]] = alloca [3 x i32]
   // B64: %[[BLOCK_SIZES6:.*]] = alloca [3 x i64]
   // CHECK-LIFETIMES: %[[BLOCK_SIZES6:.*]] = alloca [3 x i64]
   // B32: %[[BLOCK_SIZES7:.*]] = alloca [1 x i32]
   // B64: %[[BLOCK_SIZES7:.*]] = alloca [1 x i64]
   // CHECK-LIFETIMES: %[[BLOCK_SIZES7:.*]] = alloca [1 x i64]
 
   // Emits block literal on stack and block kernel [[INVLK1]].
-  // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+  // SPIR: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
+  // X86: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
   // COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
   // COMMON: store ptr addrspace(4) addrspacecast (ptr [[INVL1:@__device_side_enqueue_block_invoke[^ ]*]] to ptr addrspace(4)), ptr %block.invoke
   // COMMON: [[BL_I8:%[0-9]+]] ={{.*}} addrspacecast ptr %block to ptr addrspace(4)
-  // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_basic(
-  // COMMON-SAME: ptr [[DEF_Q]], i32 [[FLAGS]], ptr byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
+  // COMMON-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_basic(
+  // SPIR-SAME: target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], ptr byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
+  // X86-SAME: ptr [[DEF_Q]], i32 [[FLAGS]], ptr byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVLK1:[^ ]+_kernel]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) [[BL_I8]])
   enqueue_kernel(default_queue, flags, ndrange,
                  ^(void) {
                    a[i] = b[i];
                  });
 
   // Emits block literal on stack and block kernel [[INVLK2]].
-  // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+  // SPIR: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
+  // X86: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
   // COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
   // COMMON: [[WAIT_EVNT:%[0-9]+]] ={{.*}} addrspacecast ptr %event_wait_list to ptr addrspace(4)
   // COMMON: [[EVNT:%[0-9]+]] ={{.*}} addrspacecast ptr %clk_event to ptr addrspace(4)
   // COMMON: store ptr addrspace(4) addrspacecast (ptr [[INVL2:@__device_side_enqueue_block_invoke[^ ]*]] to ptr addrspace(4)), ptr %block.invoke
   // COMMON: [[BL_I8:%[0-9]+]] ={{.*}} addrspacecast ptr %block4 to ptr addrspace(4)
-  // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_basic_events
-  // COMMON-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]],  ptr {{.*}}, i32 2, ptr addrspace(4) [[WAIT_EVNT]], ptr addrspace(4) [[EVNT]],
+  // COMMON-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_basic_events
+  // SPIR-SAME: (target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]],  ptr {{.*}}, i32 2, ptr addrspace(4) [[WAIT_EVNT]], ptr addrspace(4) [[EVNT]],
+  // X86-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]],  ptr {{.*}}, i32 2, ptr addrspace(4) [[WAIT_EVNT]], ptr addrspace(4) [[EVNT]],
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVLK2:[^ ]+_kernel]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) [[BL_I8]])
   enqueue_kernel(default_queue, flags, ndrange, 2, &event_wait_list, &clk_event,
                  ^(void) {
                    a[i] = b[i];
                  });
 
-  // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_basic_events
-  // COMMON-SAME: (ptr {{%[0-9]+}}, i32 {{%[0-9]+}}, ptr {{.*}}, i32 1, ptr addrspace(4) null, ptr addrspace(4) null,
+  // COMMON-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_basic_events
+  // SPIR-SAME: (target("spirv.Queue") {{%[0-9]+}}, i32 {{%[0-9]+}}, ptr {{.*}}, i32 1, ptr addrspace(4) null, ptr addrspace(4) null,
+  // X86-SAME: (ptr {{%[0-9]+}}, i32 {{%[0-9]+}}, ptr {{.*}}, i32 1, ptr addrspace(4) null, ptr addrspace(4) null,
   enqueue_kernel(default_queue, flags, ndrange, 1, 0, 0,
                  ^(void) {
                    return;
                  });
 
   // Emits global block literal [[BLG1]] and block kernel [[INVGK1]].
-  // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+  // SPIR: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
+  // X86: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
   // COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
   // CHECK-LIFETIMES: call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %[[BLOCK_SIZES1]])
-  // CHECK-LIFETIMES-LABEL: call spir_func i32 @__enqueue_kernel_varargs(
+  // CHECK-LIFETIMES-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_varargs(
   // CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %[[BLOCK_SIZES1]])
   // B32: %[[TMP:.*]] = getelementptr [1 x i32], ptr %[[BLOCK_SIZES1]], i32 0, i32 0
   // B32: store i32 256, ptr %[[TMP]], align 4
   // B64: %[[TMP:.*]] = getelementptr [1 x i64], ptr %[[BLOCK_SIZES1]], i32 0, i32 0
   // B64: store i64 256, ptr %[[TMP]], align 8
-  // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_varargs(
-  // COMMON-SAME: ptr [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
+  // COMMON-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_varargs(
+  // SPIR-SAME: target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
+  // X86-SAME: ptr [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK1:[^ ]+_kernel]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG1]] to ptr addrspace(4)), i32 1,
   // B32-SAME: ptr %[[TMP]])
   // B64-SAME: ptr %[[TMP]])
   enqueue_kernel(default_queue, flags, ndrange,
                  ^(local void *p) {
                    return;
                  },
                  256);
 
   char c;
   // Emits global block literal [[BLG2]] and block kernel [[INVGK2]].
-  // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+  // SPIR: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
+  // X86: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
   // COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
   // CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %[[BLOCK_SIZES2]])
-  // CHECK-LIFETIMES-LABEL: call spir_func i32 @__enqueue_kernel_varargs(
+  // CHECK-LIFETIMES-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_varargs(
   // CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %[[BLOCK_SIZES2]])
   // B32: %[[TMP:.*]] = getelementptr [1 x i32], ptr %[[BLOCK_SIZES2]], i32 0, i32 0
   // B32: store i32 %{{.*}}, ptr %[[TMP]], align 4
   // B64: %[[TMP:.*]] = getelementptr [1 x i64], ptr %[[BLOCK_SIZES2]], i32 0, i32 0
   // B64: store i64 %{{.*}}, ptr %[[TMP]], align 8
-  // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_varargs(
-  // COMMON-SAME: ptr [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
+  // COMMON-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_varargs(
+  // SPIR-SAME: target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
+  // X86-SAME: ptr [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK2:[^ ]+_kernel]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG2]] to ptr addrspace(4)), i32 1,
   // B32-SAME: ptr %[[TMP]])
   // B64-SAME: ptr %[[TMP]])
   enqueue_kernel(default_queue, flags, ndrange,
                  ^(local void *p) {
                    return;
                  },
                  c);
 
   // Emits global block literal [[BLG3]] and block kernel [[INVGK3]].
-  // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+  // SPIR: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
+  // X86: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
   // COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
-  // COMMON: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x ptr], ptr %event_wait_list2, i{{32|64}} 0, i{{32|64}} 0
+  // SPIR: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x target("spirv.DeviceEvent")], ptr %event_wait_list2, i{{32|64}} 0, i{{32|64}} 0
+  // X86: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x ptr], ptr %event_wait_list2, i{{32|64}} 0, i{{32|64}} 0
   // COMMON: [[WAIT_EVNT:%[0-9]+]] ={{.*}} addrspacecast ptr [[AD]] to ptr addrspace(4)
   // COMMON: [[EVNT:%[0-9]+]]  ={{.*}} addrspacecast ptr %clk_event to ptr addrspace(4)
   // CHECK-LIFETIMES: call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %[[BLOCK_SIZES3]])
-  // CHECK-LIFETIMES-LABEL: call spir_func i32 @__enqueue_kernel_events_varargs(
+  // CHECK-LIFETIMES-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_events_varargs(
   // CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %[[BLOCK_SIZES3]])
   // B32: %[[TMP:.*]] = getelementptr [1 x i32], ptr %[[BLOCK_SIZES3]], i32 0, i32 0
   // B32: store i32 256, ptr %[[TMP]], align 4
   // B64: %[[TMP:.*]] = getelementptr [1 x i64], ptr %[[BLOCK_SIZES3]], i32 0, i32 0
   // B64: store i64 256, ptr %[[TMP]], align 8
-  // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_events_varargs
-  // COMMON-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]],  ptr {{.*}}, i32 2, ptr addrspace(4) [[WAIT_EVNT]], ptr addrspace(4) [[EVNT]],
+  // COMMON-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_events_varargs
+  // SPIR-SAME: (target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]],  ptr {{.*}}, i32 2, ptr addrspace(4) [[WAIT_EVNT]], ptr addrspace(4) [[EVNT]],
+  // X86-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]],  ptr {{.*}}, i32 2, ptr addrspace(4) [[WAIT_EVNT]], ptr addrspace(4) [[EVNT]],
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK3:[^ ]+_kernel]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG3]] to ptr addrspace(4)), i32 1,
   // B32-SAME: ptr %[[TMP]])
   // B64-SAME: ptr %[[TMP]])
   enqueue_kernel(default_queue, flags, ndrange, 2, event_wait_list2, &clk_event,
                  ^(local void *p) {
                    return;
                  },
                  256);
 
   // Emits global block literal [[BLG4]] and block kernel [[INVGK4]].
-  // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+  // SPIR: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
+  // X86: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
   // COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
-  // COMMON: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x ptr], ptr %event_wait_list2, i{{32|64}} 0, i{{32|64}} 0
+  // SPIR: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x target("spirv.DeviceEvent")], ptr %event_wait_list2, i{{32|64}} 0, i{{32|64}} 0
+  // X86: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x ptr], ptr %event_wait_list2, i{{32|64}} 0, i{{32|64}} 0
   // COMMON: [[WAIT_EVNT:%[0-9]+]] ={{.*}} addrspacecast ptr [[AD]] to ptr addrspace(4)
   // COMMON: [[EVNT:%[0-9]+]]  ={{.*}} addrspacecast ptr %clk_event to ptr addrspace(4)
   // CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %[[BLOCK_SIZES4]])
-  // CHECK-LIFETIMES-LABEL: call spir_func i32 @__enqueue_kernel_events_varargs(
+  // CHECK-LIFETIMES-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_events_varargs(
   // CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %[[BLOCK_SIZES4]])
   // B32: %[[TMP:.*]] = getelementptr [1 x i32], ptr %[[BLOCK_SIZES4]], i32 0, i32 0
   // B32: store i32 %{{.*}}, ptr %[[TMP]], align 4
   // B64: %[[TMP:.*]] = getelementptr [1 x i64], ptr %[[BLOCK_SIZES4]], i32 0, i32 0
   // B64: store i64 %{{.*}}, ptr %[[TMP]], align 8
-  // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_events_varargs
-  // COMMON-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]],  ptr {{.*}}, i32 2, ptr addrspace(4) [[WAIT_EVNT]], ptr addrspace(4) [[EVNT]],
+  // COMMON-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_events_varargs
+  // SPIR-SAME: (target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]],  ptr {{.*}}, i32 2, ptr addrspace(4) [[WAIT_EVNT]], ptr addrspace(4) [[EVNT]],
+  // X86-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]],  ptr {{.*}}, i32 2, ptr addrspace(4) [[WAIT_EVNT]], ptr addrspace(4) [[EVNT]],
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK4:[^ ]+_kernel]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG4]] to ptr addrspace(4)), i32 1,
   // B32-SAME: ptr %[[TMP]])
   // B64-SAME: ptr %[[TMP]])
   enqueue_kernel(default_queue, flags, ndrange, 2, event_wait_list2, &clk_event,
                  ^(local void *p) {
                    return;
                  },
                  c);
 
   long l;
   // Emits global block literal [[BLG5]] and block kernel [[INVGK5]].
-  // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+  // SPIR: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
+  // X86: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
   // COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
   // CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %[[BLOCK_SIZES5]])
-  // CHECK-LIFETIMES-LABEL: call spir_func i32 @__enqueue_kernel_varargs(
+  // CHECK-LIFETIMES-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_varargs(
   // CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %[[BLOCK_SIZES5]])
   // B32: %[[TMP:.*]] = getelementptr [1 x i32], ptr %[[BLOCK_SIZES5]], i32 0, i32 0
   // B32: store i32 %{{.*}}, ptr %[[TMP]], align 4
   // B64: %[[TMP:.*]] = getelementptr [1 x i64], ptr %[[BLOCK_SIZES5]], i32 0, i32 0
   // B64: store i64 %{{.*}}, ptr %[[TMP]], align 8
-  // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_varargs
-  // COMMON-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
+  // COMMON-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_varargs
+  // SPIR-SAME: (target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
+  // X86-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK5:[^ ]+_kernel]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG5]] to ptr addrspace(4)), i32 1,
   // B32-SAME: ptr %[[TMP]])
   // B64-SAME: ptr %[[TMP]])
   enqueue_kernel(default_queue, flags, ndrange,
                  ^(local void *p) {
                    return;
                  },
                  l);
 
   // Emits global block literal [[BLG6]] and block kernel [[INVGK6]].
-  // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+  // SPIR: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
+  // X86: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
   // COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
   // CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.start.p0(i64 24, ptr nonnull %[[BLOCK_SIZES6]])
-  // CHECK-LIFETIMES-LABEL: call spir_func i32 @__enqueue_kernel_varargs(
+  // CHECK-LIFETIMES-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_varargs(
   // CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.end.p0(i64 24, ptr nonnull %[[BLOCK_SIZES6]])
   // B32: %[[TMP:.*]] = getelementptr [3 x i32], ptr %[[BLOCK_SIZES6]], i32 0, i32 0
   // B32: store i32 1, ptr %[[TMP]], align 4
   // B32: %[[BLOCK_SIZES62:.*]] = getelementptr [3 x i32], ptr %[[BLOCK_SIZES6]], i32 0, i32 1
   // B32: store i32 2, ptr %[[BLOCK_SIZES62]], align 4
   // B32: %[[BLOCK_SIZES63:.*]] = getelementptr [3 x i32], ptr %[[BLOCK_SIZES6]], i32 0, i32 2
   // B32: store i32 4, ptr %[[BLOCK_SIZES63]], align 4
   // B64: %[[TMP:.*]] = getelementptr [3 x i64], ptr %[[BLOCK_SIZES6]], i32 0, i32 0
   // B64: store i64 1, ptr %[[TMP]], align 8
   // B64: %[[BLOCK_SIZES62:.*]] = getelementptr [3 x i64], ptr %[[BLOCK_SIZES6]], i32 0, i32 1
   // B64: store i64 2, ptr %[[BLOCK_SIZES62]], align 8
   // B64: %[[BLOCK_SIZES63:.*]] = getelementptr [3 x i64], ptr %[[BLOCK_SIZES6]], i32 0, i32 2
   // B64: store i64 4, ptr %[[BLOCK_SIZES63]], align 8
-  // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_varargs
-  // COMMON-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
+  // COMMON-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_varargs
+  // SPIR-SAME: (target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
+  // X86-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK6:[^ ]+_kernel]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG6]] to ptr addrspace(4)), i32 3,
   // B32-SAME: ptr %[[TMP]])
   // B64-SAME: ptr %[[TMP]])
   enqueue_kernel(default_queue, flags, ndrange,
                  ^(local void *p1, local void *p2, local void *p3) {
                    return;
                  },
                  1, 2, 4);
 
   // Emits global block literal [[BLG7]] and block kernel [[INVGK7]].
-  // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+  // SPIR: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
+  // X86: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
   // COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
   // CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %[[BLOCK_SIZES7]])
-  // CHECK-LIFETIMES-LABEL: call spir_func i32 @__enqueue_kernel_varargs(
+  // CHECK-LIFETIMES-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_varargs(
   // CHECK-LIFETIMES-NEXT: call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %[[BLOCK_SIZES7]])
   // B32: %[[TMP:.*]] = getelementptr [1 x i32], ptr %[[BLOCK_SIZES7]], i32 0, i32 0
   // B32: store i32 0, ptr %[[TMP]], align 4
   // B64: %[[TMP:.*]] = getelementptr [1 x i64], ptr %[[BLOCK_SIZES7]], i32 0, i32 0
   // B64: store i64 4294967296, ptr %[[TMP]], align 8
-  // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_varargs
-  // COMMON-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
+  // COMMON-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_varargs
+  // SPIR-SAME: (target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
+  // X86-SAME: (ptr [[DEF_Q]], i32 [[FLAGS]], ptr [[NDR]]{{([0-9]+)?}},
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK7:[^ ]+_kernel]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG7]] to ptr addrspace(4)), i32 1,
   // B32-SAME: ptr %[[TMP]])
   // B64-SAME: ptr %[[TMP]])
   enqueue_kernel(default_queue, flags, ndrange,
                  ^(local void *p) {
                    return;
                  },
kernel void device_side_enqueue(global int *a, global int *b, int i) {
 
   // Emits global block literal [[BLG9]] and invoke function [[INVG9]].
   // COMMON: store ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG9]] to ptr addrspace(4)), ptr %block_B
   void (^const block_B)(local void *) = ^(local void *a) {
     return;
   };
 
   // Uses global block literal [[BLG8]] and invoke function [[INVG8]].
-  // COMMON: call spir_func void @__device_side_enqueue_block_invoke_11(ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG8]] to ptr addrspace(4))) [[INVOKE_ATTR:#[0-9]+]]
+  // COMMON: call {{(spir_func )?}}void @__device_side_enqueue_block_invoke_11(ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG8]] to ptr addrspace(4))) [[INVOKE_ATTR:#[0-9]+]]
   block_A();
 
   // Emits global block literal [[BLG8]] and block kernel [[INVGK8]]. [[INVGK8]] calls [[INVG8]].
-  // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+  // SPIR: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
+  // X86: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
   // COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
-  // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_basic(
-  // COMMON-SAME: ptr [[DEF_Q]], i32 [[FLAGS]], ptr byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
+  // COMMON-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_basic(
+  // SPIR-SAME: target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], ptr byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
+  // X86-SAME: ptr [[DEF_Q]], i32 [[FLAGS]], ptr byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK8:[^ ]+_kernel]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG8]] to ptr addrspace(4)))
   enqueue_kernel(default_queue, flags, ndrange, block_A);
 
   // Uses block kernel [[INVGK8]] and global block literal [[BLG8]].
-  // COMMON: call spir_func i32 @__get_kernel_work_group_size_impl(
+  // COMMON: call {{(spir_func )?}}i32 @__get_kernel_work_group_size_impl(
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK8]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG8]] to ptr addrspace(4)))
   unsigned size = get_kernel_work_group_size(block_A);
 
   // Uses global block literal [[BLG8]] and invoke function [[INVG8]]. Make sure no redundant block literal and invoke functions are emitted.
-  // COMMON: call spir_func void @__device_side_enqueue_block_invoke_11(ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG8]] to ptr addrspace(4)))
+  // COMMON: call {{(spir_func )?}}void @__device_side_enqueue_block_invoke_11(ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG8]] to ptr addrspace(4)))
   block_A();
 
   // Make sure that block invoke function is resolved correctly after sequence of assignements.
   // COMMON: store ptr addrspace(4)
   // COMMON-SAME: addrspacecast (ptr addrspace(1)
   // COMMON-SAME: [[BL_GLOBAL]]
   // COMMON-SAME: to ptr addrspace(4)),
   // COMMON-SAME: ptr %b1,
   bl_t b1 = block_G;
   // COMMON: store ptr addrspace(4)
   // COMMON-SAME: addrspacecast (ptr addrspace(1)
   // COMMON-SAME: [[BL_GLOBAL]]
   // COMMON-SAME: to ptr addrspace(4)),
   // COMMON-SAME: ptr %b2,
   bl_t b2 = b1;
-  // COMMON: call spir_func void @block_G_block_invoke(ptr addrspace(4) addrspacecast (ptr addrspace(1)
+  // COMMON: call {{(spir_func )?}}void @block_G_block_invoke(ptr addrspace(4) addrspacecast (ptr addrspace(1)
   // COMMON-SAME: [[BL_GLOBAL]]
   // COOMON-SAME: to ptr addrspace(4)), ptr addrspace(3) null)
   b2(0);
   // Uses global block literal [[BL_GLOBAL]] and block kernel [[INV_G_K]]. [[INV_G_K]] calls [[INV_G]].
-  // COMMON: call spir_func i32 @__get_kernel_preferred_work_group_size_multiple_impl(
+  // COMMON: call {{(spir_func )?}}i32 @__get_kernel_preferred_work_group_size_multiple_impl(
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INV_G_K:[^ ]+_kernel]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BL_GLOBAL]] to ptr addrspace(4)))
   size = get_kernel_preferred_work_group_size_multiple(b2);
 
   void (^block_C)(void) = ^{
     callee(i, a);
   };
   // Emits block literal on stack and block kernel [[INVLK3]].
   // COMMON: store ptr addrspace(4) addrspacecast (ptr [[INVL3:@__device_side_enqueue_block_invoke[^ ]*]] to ptr addrspace(4)), ptr %block.invoke
-  // COMMON: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
+  // SPIR: [[DEF_Q:%[0-9]+]] = load target("spirv.Queue"), ptr %default_queue
+  // X86: [[DEF_Q:%[0-9]+]] = load ptr, ptr %default_queue
   // COMMON: [[FLAGS:%[0-9]+]] = load i32, ptr %flags
   // COMMON: [[BL_I8:%[0-9]+]] ={{.*}} addrspacecast ptr {{.*}} to ptr addrspace(4)
-  // COMMON-LABEL: call spir_func i32 @__enqueue_kernel_basic(
-  // COMMON-SAME: ptr [[DEF_Q]], i32 [[FLAGS]], ptr byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
+  // COMMON-LABEL: call {{(spir_func )?}}i32 @__enqueue_kernel_basic(
+  // SPIR-SAME: target("spirv.Queue") [[DEF_Q]], i32 [[FLAGS]], ptr byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
+  // X86-SAME: ptr [[DEF_Q]], i32 [[FLAGS]], ptr byval(%struct.ndrange_t) [[NDR]]{{([0-9]+)?}},
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVLK3:[^ ]+_kernel]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) [[BL_I8]])
   enqueue_kernel(default_queue, flags, ndrange, block_C);
 
   // Emits global block literal [[BLG9]] and block kernel [[INVGK9]]. [[INVGK9]] calls [[INV9]].
-  // COMMON: call spir_func i32 @__get_kernel_work_group_size_impl(
+  // COMMON: call {{(spir_func )?}}i32 @__get_kernel_work_group_size_impl(
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK9:[^ ]+_kernel]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG9]] to ptr addrspace(4)))
   size = get_kernel_work_group_size(block_B);
 
   // Uses global block literal [[BLG8]] and block kernel [[INVGK8]]. Make sure no redundant block literal ind invoke functions are emitted.
-  // COMMON: call spir_func i32 @__get_kernel_preferred_work_group_size_multiple_impl(
+  // COMMON: call {{(spir_func )?}}i32 @__get_kernel_preferred_work_group_size_multiple_impl(
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK8]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG8]] to ptr addrspace(4)))
   size = get_kernel_preferred_work_group_size_multiple(block_A);
 
   // Uses global block literal [[BL_GLOBAL]] and block kernel [[INV_G_K]]. [[INV_G_K]] calls [[INV_G]].
-  // COMMON: call spir_func i32 @__get_kernel_preferred_work_group_size_multiple_impl(
+  // COMMON: call {{(spir_func )?}}i32 @__get_kernel_preferred_work_group_size_multiple_impl(
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INV_G_K:[^ ]+_kernel]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BL_GLOBAL]] to ptr addrspace(4)))
   size = get_kernel_preferred_work_group_size_multiple(block_G);
 
   // Emits global block literal [[BLG10]] and block kernel [[INVGK10]].
-  // COMMON: call spir_func i32 @__get_kernel_max_sub_group_size_for_ndrange_impl(ptr {{[^,]+}},
+  // COMMON: call {{(spir_func )?}}i32 @__get_kernel_max_sub_group_size_for_ndrange_impl(ptr {{[^,]+}},
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK10:[^ ]+_kernel]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG10]] to ptr addrspace(4)))
   size = get_kernel_max_sub_group_size_for_ndrange(ndrange, ^(){});
 
   // Emits global block literal [[BLG11]] and block kernel [[INVGK11]].
-  // COMMON: call spir_func i32 @__get_kernel_sub_group_count_for_ndrange_impl(ptr {{[^,]+}},
+  // COMMON: call {{(spir_func )?}}i32 @__get_kernel_sub_group_count_for_ndrange_impl(ptr {{[^,]+}},
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr [[INVGK11:[^ ]+_kernel]] to ptr addrspace(4)),
   // COMMON-SAME: ptr addrspace(4) addrspacecast (ptr addrspace(1) [[BLG11]] to ptr addrspace(4)))
   size = get_kernel_sub_group_count_for_ndrange(ndrange, ^(){});
 }
 
 // COMMON: define spir_kernel void [[INVLK1]](ptr addrspace(4) %0) #{{[0-9]+}} {
 // COMMON: entry:
-// COMMON:  call spir_func void @__device_side_enqueue_block_invoke(ptr addrspace(4) %0)
+// COMMON:  call {{(spir_func )?}}void @__device_side_enqueue_block_invoke(ptr addrspace(4) %0)
 // COMMON:  ret void
 // COMMON: }
 // COMMON: define spir_kernel void [[INVLK2]](ptr addrspace(4){{.*}})
 // COMMON: define spir_kernel void [[INVGK1]](ptr addrspace(4){{.*}}, ptr addrspace(3){{.*}})  [[INVOKE_KERNEL_ATTR:#[0-9]+]]
 // COMMON: define spir_kernel void [[INVGK2]](ptr addrspace(4){{.*}}, ptr addrspace(3){{.*}})
 // COMMON: define spir_kernel void [[INVGK3]](ptr addrspace(4){{.*}}, ptr addrspace(3){{.*}})
 // COMMON: define spir_kernel void [[INVGK4]](ptr addrspace(4){{.*}}, ptr addrspace(3){{.*}})
 // COMMON: define spir_kernel void [[INVGK5]](ptr addrspace(4){{.*}}, ptr addrspace(3){{.*}})
 // COMMON: define spir_kernel void [[INVGK6]](ptr addrspace(4) %0, ptr addrspace(3) %1, ptr addrspace(3) %2, ptr addrspace(3) %3) #{{[0-9]+}} {
 // COMMON: entry:
-// COMMON:  call spir_func void @__device_side_enqueue_block_invoke_9(ptr addrspace(4) %0, ptr addrspace(3) %1, ptr addrspace(3) %2, ptr addrspace(3) %3)
+// COMMON:  call {{(spir_func )?}}void @__device_side_enqueue_block_invoke_9(ptr addrspace(4) %0, ptr addrspace(3) %1, ptr addrspace(3) %2, ptr addrspace(3) %3)
 // COMMON:  ret void
 // COMMON: }
 // COMMON: define spir_kernel void [[INVGK7]](ptr addrspace(4){{.*}}, ptr addrspace(3){{.*}})
-// COMMON: define internal spir_func void [[INVG8]](ptr addrspace(4){{.*}}) [[INVG8_INVOKE_FUNC_ATTR:#[0-9]+]]
-// COMMON: define internal spir_func void [[INVG9]](ptr addrspace(4){{.*}}, ptr addrspace(3) %{{.*}})
+// COMMON: define internal {{(spir_func )?}}void [[INVG8]](ptr addrspace(4){{.*}}) [[INVG8_INVOKE_FUNC_ATTR:#[0-9]+]]
+// COMMON: define internal {{(spir_func )?}}void [[INVG9]](ptr addrspace(4){{.*}}, ptr addrspace(3) %{{.*}})
 // COMMON: define spir_kernel void [[INVGK8]](ptr addrspace(4){{.*}})
 // COMMON: define spir_kernel void [[INV_G_K]](ptr addrspace(4){{.*}}, ptr addrspace(3){{.*}})
 // COMMON: define spir_kernel void [[INVLK3]](ptr addrspace(4){{.*}})
 // COMMON: define spir_kernel void [[INVGK9]](ptr addrspace(4){{.*}}, ptr addrspace(3){{.*}})
 // COMMON: define spir_kernel void [[INVGK10]](ptr addrspace(4){{.*}})
 // COMMON: define spir_kernel void [[INVGK11]](ptr addrspace(4){{.*}})
 
-// COMMON: attributes [[INVG8_INVOKE_FUNC_ATTR]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
-// COMMON: attributes [[INVOKE_KERNEL_ATTR]] = { convergent nounwind "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// SPIR: attributes [[INVG8_INVOKE_FUNC_ATTR]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// SPIR: attributes [[INVOKE_KERNEL_ATTR]] = { convergent nounwind "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// X86: attributes [[INVG8_INVOKE_FUNC_ATTR]] = { convergent noinline nounwind optnone "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="{{[^"]*}}" }
+// X86: attributes [[INVOKE_KERNEL_ATTR]] = { convergent nounwind "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="{{[^"]*}}" }

clang/test/CodeGenOpenCL/intel-subgroups-avc-ext-types.cl

-// RUN: %clang_cc1 -no-opaque-pointers %s -triple spir-unknown-unknown -cl-std=CL1.2 -cl-ext=+cl_intel_device_side_avc_motion_estimation -emit-llvm -o - -O0 | FileCheck %s
+// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=CL1.2 -cl-ext=+cl_intel_device_side_avc_motion_estimation -emit-llvm -o - -O0 | FileCheck %s
 
-// CHECK: %opencl.intel_sub_group_avc_mce_payload_t = type opaque
-// CHECK: %opencl.intel_sub_group_avc_ime_payload_t = type opaque
-// CHECK: %opencl.intel_sub_group_avc_ref_payload_t = type opaque
-// CHECK: %opencl.intel_sub_group_avc_sic_payload_t = type opaque
-
-// CHECK: %opencl.intel_sub_group_avc_mce_result_t = type opaque
-// CHECK: %opencl.intel_sub_group_avc_ime_result_t = type opaque
-// CHECK: %opencl.intel_sub_group_avc_ref_result_t = type opaque
-// CHECK: %opencl.intel_sub_group_avc_sic_result_t = type opaque
-
-// CHECK: %opencl.intel_sub_group_avc_ime_result_single_reference_streamout_t = type opaque
-// CHECK: %opencl.intel_sub_group_avc_ime_result_dual_reference_streamout_t = type opaque
-// CHECK: %opencl.intel_sub_group_avc_ime_single_reference_streamin_t = type opaque
-// CHECK: %opencl.intel_sub_group_avc_ime_dual_reference_streamin_t = type opaque
-
-// CHECK: store %opencl.intel_sub_group_avc_ime_payload_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ref_payload_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_sic_payload_t* null,
-
-// CHECK: store %opencl.intel_sub_group_avc_ime_result_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ref_result_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_sic_result_t* null,
-
-// CHECK: store %opencl.intel_sub_group_avc_ime_result_single_reference_streamout_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ime_result_dual_reference_streamout_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ime_single_reference_streamin_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ime_dual_reference_streamin_t* null,
-//
-// CHECK: store %opencl.intel_sub_group_avc_ime_payload_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ref_payload_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_sic_payload_t* null,
-
-// CHECK: store %opencl.intel_sub_group_avc_ime_result_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ref_result_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_sic_result_t* null,
-
-// CHECK: store %opencl.intel_sub_group_avc_ime_result_single_reference_streamout_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ime_result_dual_reference_streamout_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ime_single_reference_streamin_t* null,
-// CHECK: store %opencl.intel_sub_group_avc_ime_dual_reference_streamin_t* null,
+// CHECK: store target("spirv.AvcImePayloadINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcRefPayloadINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcSicPayloadINTEL") zeroinitializer,
+
+// CHECK: store target("spirv.AvcImeResultINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcRefResultINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcSicResultINTEL") zeroinitializer,
+
+// CHECK: store target("spirv.AvcImeResultSingleReferenceStreamoutINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcImeResultDualReferenceStreamoutINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcImeSingleReferenceStreaminINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcImeDualReferenceStreaminINTEL") zeroinitializer,
+
+// CHECK: store target("spirv.AvcImePayloadINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcRefPayloadINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcSicPayloadINTEL") zeroinitializer,
+
+// CHECK: store target("spirv.AvcImeResultINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcRefResultINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcSicResultINTEL") zeroinitializer,
+
+// CHECK: store target("spirv.AvcImeResultSingleReferenceStreamoutINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcImeResultDualReferenceStreamoutINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcImeSingleReferenceStreaminINTEL") zeroinitializer,
+// CHECK: store target("spirv.AvcImeDualReferenceStreaminINTEL") zeroinitializer,
 
 #pragma OPENCL EXTENSION cl_intel_device_side_avc_motion_estimation : enable
 
 // Using 0x0 directly allows us not to include opencl-c.h header and not to
 // redefine all of these CLK_AVC_*_INTITIALIZE_INTEL macro. '0x0' value must
 // be in sync with ones defined in opencl-c.h
 
 void foo() {

clang/test/CodeGenOpenCL/opencl_types.cl

 // RUN: %clang_cc1 -cl-std=CL2.0 %s -triple "spir-unknown-unknown" -emit-llvm -o - -O0 | FileCheck %s --check-prefixes=CHECK-COM,CHECK-SPIR
 // RUN: %clang_cc1 -cl-std=CL2.0 %s -triple "amdgcn--amdhsa" -emit-llvm -o - -O0 | FileCheck %s --check-prefixes=CHECK-COM,CHECK-AMDGCN
 
 #define CLK_ADDRESS_CLAMP_TO_EDGE       2
 #define CLK_NORMALIZED_COORDS_TRUE      1
 #define CLK_FILTER_NEAREST              0x10
 #define CLK_FILTER_LINEAR               0x20
 
 constant sampler_t glb_smp = CLK_ADDRESS_CLAMP_TO_EDGE|CLK_NORMALIZED_COORDS_TRUE|CLK_FILTER_NEAREST;
 // CHECK-COM-NOT: constant i32
 
 void fnc1(image1d_t img) {}
-// CHECK-SPIR: @fnc1(ptr addrspace(1)
+// CHECK-SPIR: @fnc1(target("spirv.Image", void, 0, 0, 0, 0, 0, 0, 0)
 // CHECK-AMDGCN: @fnc1(ptr addrspace(4)
 
 void fnc1arr(image1d_array_t img) {}
-// CHECK-SPIR: @fnc1arr(ptr addrspace(1)
+// CHECK-SPIR: @fnc1arr(target("spirv.Image", void, 0, 0, 1, 0, 0, 0, 0)
 // CHECK-AMDGCN: @fnc1arr(ptr addrspace(4)
 
 void fnc1buff(image1d_buffer_t img) {}
-// CHECK-SPIR: @fnc1buff(ptr addrspace(1)
+// CHECK-SPIR: @fnc1buff(target("spirv.Image", void, 5, 0, 0, 0, 0, 0, 0)
 // CHECK-AMDGCN: @fnc1buff(ptr addrspace(4)
 
 void fnc2(image2d_t img) {}
-// CHECK-SPIR: @fnc2(ptr addrspace(1)
+// CHECK-SPIR: @fnc2(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0)
 // CHECK-AMDGCN: @fnc2(ptr addrspace(4)
 
 void fnc2arr(image2d_array_t img) {}
-// CHECK-SPIR: @fnc2arr(ptr addrspace(1)
+// CHECK-SPIR: @fnc2arr(target("spirv.Image", void, 1, 0, 1, 0, 0, 0, 0)
 // CHECK-AMDGCN: @fnc2arr(ptr addrspace(4)
 
 void fnc3(image3d_t img) {}
-// CHECK-SPIR: @fnc3(ptr addrspace(1)
+// CHECK-SPIR: @fnc3(target("spirv.Image", void, 2, 0, 0, 0, 0, 0, 0)
 // CHECK-AMDGCN: @fnc3(ptr addrspace(4)
 
 void fnc4smp(sampler_t s) {}
-// CHECK-SPIR-LABEL: define {{.*}}void @fnc4smp(ptr addrspace(2)
+// CHECK-SPIR-LABEL: define {{.*}}void @fnc4smp(target("spirv.Sampler")
 // CHECK-AMDGCN-LABEL: define {{.*}}void @fnc4smp(ptr addrspace(4)
 
 kernel void foo(image1d_t img) {
   sampler_t smp = CLK_ADDRESS_CLAMP_TO_EDGE|CLK_NORMALIZED_COORDS_TRUE|CLK_FILTER_LINEAR;
-  // CHECK-SPIR: alloca ptr addrspace(2)
+  // CHECK-SPIR: alloca target("spirv.Sampler")
   // CHECK-AMDGCN: alloca ptr addrspace(4)
   event_t evt;
-  // CHECK-SPIR: alloca ptr
+  // CHECK-SPIR: alloca target("spirv.Event")
   // CHECK-AMDGCN: alloca ptr addrspace(5)
   clk_event_t clk_evt;
-  // CHECK-SPIR: alloca ptr
+  // CHECK-SPIR: alloca target("spirv.DeviceEvent")
   // CHECK-AMDGCN: alloca ptr addrspace(1)
   queue_t queue;
-  // CHECK-SPIR: alloca ptr
+  // CHECK-SPIR: alloca target("spirv.Queue")
   // CHECK-AMDGCN: alloca ptr addrspace(1)
   reserve_id_t rid;
-  // CHECK-SPIR: alloca ptr
+  // CHECK-SPIR: alloca target("spirv.ReserveId")
   // CHECK-AMDGCN: alloca ptr addrspace(1)
-  // CHECK-SPIR: store ptr addrspace(2)
+  // CHECK-SPIR: store target("spirv.Sampler")
   // CHECK-AMDGCN: store ptr addrspace(4)
   fnc4smp(smp);
-  // CHECK-SPIR: call {{.*}}void @fnc4smp(ptr addrspace(2)
+  // CHECK-SPIR: call {{.*}}void @fnc4smp(target("spirv.Sampler")
   // CHECK-AMDGCN: call {{.*}}void @fnc4smp(ptr addrspace(4)
   fnc4smp(glb_smp);
-  // CHECK-SPIR: call {{.*}}void @fnc4smp(ptr addrspace(2)
+  // CHECK-SPIR: call {{.*}}void @fnc4smp(target("spirv.Sampler")
   // CHECK-AMDGCN: call {{.*}}void @fnc4smp(ptr addrspace(4)
 }
 
 kernel void foo_ro_pipe(read_only pipe int p) {}
-// CHECK-SPIR: @foo_ro_pipe(ptr addrspace(1) %p)
+// CHECK-SPIR: @foo_ro_pipe(target("spirv.Pipe", 0) %p)
 // CHECK_AMDGCN: @foo_ro_pipe(ptr addrspace(1) %p)
 
 kernel void foo_wo_pipe(write_only pipe int p) {}
-// CHECK-SPIR: @foo_wo_pipe(ptr addrspace(1) %p)
+// CHECK-SPIR: @foo_wo_pipe(target("spirv.Pipe", 1) %p)
 // CHECK_AMDGCN: @foo_wo_pipe(ptr addrspace(1) %p)
 
 void __attribute__((overloadable)) bad1(image1d_t b, image2d_t c, image2d_t d) {}
 // CHECK-SPIR-LABEL: @{{_Z4bad114ocl_image1d_ro14ocl_image2d_roS0_|"\\01\?bad1@@\$\$J0YAXPAUocl_image1d_ro@@PAUocl_image2d_ro@@1@Z"}}
 // CHECK-AMDGCN-LABEL: @{{_Z4bad114ocl_image1d_ro14ocl_image2d_roS0_|"\\01\?bad1@@\$\$J0YAXPAUocl_image1d_ro@@PAUocl_image2d_ro@@1@Z"}}(ptr addrspace(4){{.*}}ptr addrspace(4){{.*}}ptr addrspace(4){{.*}})

clang/test/CodeGenOpenCL/sampler.cl

-// RUN: %clang_cc1 -no-opaque-pointers %s -emit-llvm -triple spir-unknown-unknown -o - -O0 | FileCheck %s
-// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=CL2.0 -emit-llvm -triple spir-unknown-unknown -o - -O0 | FileCheck %s
-// RUN: %clang_cc1 -no-opaque-pointers %s -cl-std=clc++ -emit-llvm -triple spir-unknown-unknown -o - -O0 | FileCheck %s
+// RUN: %clang_cc1 %s -emit-llvm -triple spir-unknown-unknown -o - -O0 | FileCheck --check-prefixes=CHECK-COMMON,CHECK-SPIR %s
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -triple spir-unknown-unknown -o - -O0 | FileCheck --check-prefixes=CHECK-COMMON,CHECK-SPIR %s
+// RUN: %clang_cc1 %s -cl-std=clc++ -emit-llvm -triple spir-unknown-unknown -o - -O0 | FileCheck --check-prefixes=CHECK-COMMON,CHECK-SPIR %s
+// RUN: %clang_cc1 %s -emit-llvm -triple x86_64-unknown-linux-gnu -o - -O0 | FileCheck --check-prefixes=CHECK-COMMON,CHECK-X86 %s
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -triple x86_64-unknown-linux-gnu -o - -O0 | FileCheck --check-prefixes=CHECK-COMMON,CHECK-X86 %s
+// RUN: %clang_cc1 %s -cl-std=clc++ -emit-llvm -triple x86_64-unknown-linux-gnu -o - -O0 | FileCheck --check-prefixes=CHECK-COMMON,CHECK-X86 %s
 //
 // This test covers 5 cases of sampler initialzation:
 //   1. function argument passing
 //      1a. argument is a file-scope variable
 //      1b. argument is a function-scope variable
 //      1c. argument is one of caller function's parameters
 //   2. variable initialization
 //      2a. initializing a file-scope variable with constant addr space qualifier
 //      2b. initializing a function-scope variable
 //      2c. initializing a file-scope variable with const qualifier
 
 #define CLK_ADDRESS_CLAMP_TO_EDGE       2
 #define CLK_NORMALIZED_COORDS_TRUE      1
 #define CLK_FILTER_NEAREST              0x10
 #define CLK_FILTER_LINEAR               0x20
 
-// CHECK: %opencl.sampler_t = type opaque
-
 // Case 2a
 constant sampler_t glb_smp = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_LINEAR;
-// CHECK-NOT: glb_smp
+// CHECK-COMMON-NOT: glb_smp
 
 // Case 2c
 const sampler_t glb_smp_const = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_LINEAR;
-// CHECK-NOT: glb_smp_const
+// CHECK-COMMON-NOT: glb_smp_const
 
 int get_sampler_initializer(void);
 
 void fnc4smp(sampler_t s) {}
-// CHECK: define{{.*}} spir_func void [[FUNCNAME:@.*fnc4smp.*]](%opencl.sampler_t addrspace(2)* %
+// CHECK-SPIR: define{{.*}} spir_func void [[FUNCNAME:@.*fnc4smp.*]](target("spirv.Sampler") %
+// CHECK-X86: define{{.*}} void [[FUNCNAME:@.*fnc4smp.*]](ptr %
 
 kernel void foo(sampler_t smp_par) {
-  // CHECK-LABEL: define{{.*}} spir_kernel void @foo(%opencl.sampler_t addrspace(2)* %smp_par)
-  // CHECK: [[smp_par_ptr:%[A-Za-z0-9_\.]+]] = alloca %opencl.sampler_t addrspace(2)*
+  // CHECK-SPIR-LABEL: define{{.*}} spir_kernel void @foo(target("spirv.Sampler") %smp_par)
+  // CHECK-SPIR: [[smp_par_ptr:%[A-Za-z0-9_\.]+]] = alloca target("spirv.Sampler")
+  // CHECK-X86-LABEL: define{{.*}} spir_kernel void @foo(ptr %smp_par)
+  // CHECK-X86: [[smp_par_ptr:%[A-Za-z0-9_\.]+]] = alloca ptr
 
   // Case 2b
   sampler_t smp = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_NEAREST;
-  // CHECK: [[smp_ptr:%[A-Za-z0-9_\.]+]] = alloca %opencl.sampler_t addrspace(2)*
-  // CHECK: [[SAMP:%[0-9]+]] = call spir_func %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 19)
-  // CHECK: store %opencl.sampler_t addrspace(2)* [[SAMP]], %opencl.sampler_t addrspace(2)** [[smp_ptr]]
+  // CHECK-SPIR: [[smp_ptr:%[A-Za-z0-9_\.]+]] = alloca target("spirv.Sampler")
+  // CHECK-SPIR: [[SAMP:%[0-9]+]] = call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 19)
+  // CHECK-SPIR: store target("spirv.Sampler") [[SAMP]], ptr [[smp_ptr]]
+  // CHECK-X86: [[smp_ptr:%[A-Za-z0-9_\.]+]] = alloca ptr
+  // CHECK-X86: [[SAMP:%[0-9]+]] = call ptr @__translate_sampler_initializer(i32 19)
+  // CHECK-X86: store ptr [[SAMP]], ptr [[smp_ptr]]
 
   // Case 1b
   fnc4smp(smp);
-  // CHECK-NOT: call %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 19)
-  // CHECK: [[SAMP:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[smp_ptr]]
-  // CHECK: call spir_func void [[FUNCNAME]](%opencl.sampler_t addrspace(2)* [[SAMP]])
+  // CHECK-SPIR-NOT: call target("spirv.Sampler") @__translate_sampler_initializer(i32 19)
+  // CHECK-SPIR: [[SAMP:%[0-9]+]] = load target("spirv.Sampler"), ptr [[smp_ptr]]
+  // CHECK-SPIR: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[SAMP]])
+  // CHECK-X86-NOT: call ptr @__translate_sampler_initializer(i32 19)
+  // CHECK-X86: [[SAMP:%[0-9]+]] = load ptr, ptr [[smp_ptr]]
+  // CHECK-X86: call void [[FUNCNAME]](ptr [[SAMP]])
 
   // Case 1b
   fnc4smp(smp);
-  // CHECK-NOT: call %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 19)
-  // CHECK: [[SAMP:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[smp_ptr]]
-  // CHECK: call spir_func void [[FUNCNAME]](%opencl.sampler_t addrspace(2)* [[SAMP]])
+  // CHECK-SPIR-NOT: call target("spirv.Sampler") @__translate_sampler_initializer(i32 19)
+  // CHECK-SPIR: [[SAMP:%[0-9]+]] = load target("spirv.Sampler"), ptr [[smp_ptr]]
+  // CHECK-SPIR: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[SAMP]])
+  // CHECK-X86-NOT: call ptr @__translate_sampler_initializer(i32 19)
+  // CHECK-X86: [[SAMP:%[0-9]+]] = load ptr, ptr [[smp_ptr]]
+  // CHECK-X86: call void [[FUNCNAME]](ptr [[SAMP]])
 
   // Case 1a/2a
   fnc4smp(glb_smp);
-  // CHECK: [[SAMP:%[0-9]+]] = call spir_func %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 35)
-  // CHECK: call spir_func void [[FUNCNAME]](%opencl.sampler_t addrspace(2)* [[SAMP]])
+  // CHECK-SPIR: [[SAMP:%[0-9]+]] = call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 35)
+  // CHECK-SPIR: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[SAMP]])
+  // CHECK-X86: [[SAMP:%[0-9]+]] = call ptr @__translate_sampler_initializer(i32 35)
+  // CHECK-X86: call void [[FUNCNAME]](ptr [[SAMP]])
 
   // Case 1a/2c
   fnc4smp(glb_smp_const);
-  // CHECK: [[SAMP:%[0-9]+]] = call spir_func %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 35)
-  // CHECK: call spir_func void [[FUNCNAME]](%opencl.sampler_t addrspace(2)* [[SAMP]])
+  // CHECK-SPIR: [[SAMP:%[0-9]+]] = call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 35)
+  // CHECK-SPIR: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[SAMP]])
+  // CHECK-X86: [[SAMP:%[0-9]+]] = call ptr @__translate_sampler_initializer(i32 35)
+  // CHECK-X86: call void [[FUNCNAME]](ptr [[SAMP]])
 
   // Case 1c
   fnc4smp(smp_par);
-  // CHECK: [[SAMP:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[smp_par_ptr]]
-  // CHECK: call spir_func void [[FUNCNAME]](%opencl.sampler_t addrspace(2)* [[SAMP]])
+  // CHECK-SPIR: [[SAMP:%[0-9]+]] = load target("spirv.Sampler"), ptr [[smp_par_ptr]]
+  // CHECK-SPIR: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[SAMP]])
+  // CHECK-X86: [[SAMP:%[0-9]+]] = load ptr, ptr [[smp_par_ptr]]
+  // CHECK-X86: call void [[FUNCNAME]](ptr [[SAMP]])
 
   fnc4smp(5);
-  // CHECK: [[SAMP:%[0-9]+]] = call spir_func %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 5)
-  // CHECK: call spir_func void [[FUNCNAME]](%opencl.sampler_t addrspace(2)* [[SAMP]])
+  // CHECK-SPIR: [[SAMP:%[0-9]+]] = call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 5)
+  // CHECK-SPIR: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[SAMP]])
+  // CHECK-X86: [[SAMP:%[0-9]+]] = call ptr @__translate_sampler_initializer(i32 5)
+  // CHECK-X86: call void [[FUNCNAME]](ptr [[SAMP]])
 
   const sampler_t const_smp = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_LINEAR;
   fnc4smp(const_smp);
-   // CHECK: [[CONST_SAMP:%[0-9]+]] = call spir_func %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 35)
-  // CHECK: store %opencl.sampler_t addrspace(2)* [[CONST_SAMP]], %opencl.sampler_t addrspace(2)** [[CONST_SMP_PTR:%[a-zA-Z0-9]+]]
+  // CHECK-SPIR: [[CONST_SAMP:%[0-9]+]] = call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 35)
+  // CHECK-SPIR: store target("spirv.Sampler") [[CONST_SAMP]], ptr [[CONST_SMP_PTR:%[a-zA-Z0-9]+]]
+  // CHECK-X86: [[CONST_SAMP:%[0-9]+]] = call ptr @__translate_sampler_initializer(i32 35)
+  // CHECK-X86: store ptr [[CONST_SAMP]], ptr [[CONST_SMP_PTR:%[a-zA-Z0-9]+]]
   fnc4smp(const_smp);
-  // CHECK: [[SAMP:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[CONST_SMP_PTR]]
-  // CHECK: call spir_func void [[FUNCNAME]](%opencl.sampler_t addrspace(2)* [[SAMP]])
+  // CHECK-SPIR: [[SAMP:%[0-9]+]] = load target("spirv.Sampler"), ptr [[CONST_SMP_PTR]]
+  // CHECK-SPIR: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[SAMP]])
+  // CHECK-X86: [[SAMP:%[0-9]+]] = load ptr, ptr [[CONST_SMP_PTR]]
+  // CHECK-X86: call void [[FUNCNAME]](ptr [[SAMP]])
 
   constant sampler_t constant_smp = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_LINEAR;
   fnc4smp(constant_smp);
-  // CHECK: [[SAMP:%[0-9]+]] = call spir_func %opencl.sampler_t addrspace(2)* @__translate_sampler_initializer(i32 35)
-  // CHECK: call spir_func void [[FUNCNAME]](%opencl.sampler_t addrspace(2)* [[SAMP]])
+  // CHECK-SPIR: [[SAMP:%[0-9]+]] = call spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32 35)
+  // CHECK-SPIR: call spir_func void [[FUNCNAME]](target("spirv.Sampler") [[SAMP]])
+  // CHECK-X86: [[SAMP:%[0-9]+]] = call ptr @__translate_sampler_initializer(i32 35)
+  // CHECK-X86: call void [[FUNCNAME]](ptr [[SAMP]])
 
   // TODO: enable sampler initialization with non-constant integer.
   //const sampler_t const_smp_func_init = get_sampler_initializer();
 }

clang/test/Index/pipe-size.cl

 // RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple x86_64-unknown-linux-gnu %s -o - | FileCheck %s --check-prefix=X86
 // RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple spir-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPIR
 // RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple spir64-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPIR64
 // RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple amdgcn-amd-amdhsa %s -o - | FileCheck %s --check-prefix=AMDGCN
 __kernel void testPipe( pipe int test )
 {
     int s = sizeof(test);
     // X86: store ptr %test, ptr %test.addr, align 8
     // X86: store i32 8, ptr %s, align 4
-    // SPIR: store ptr addrspace(1) %test, ptr %test.addr, align 4
+    // SPIR: store target("spirv.Pipe", 0) %test, ptr %test.addr, align 4
     // SPIR: store i32 4, ptr %s, align 4
-    // SPIR64: store ptr addrspace(1) %test, ptr %test.addr, align 8
+    // SPIR64: store target("spirv.Pipe", 0) %test, ptr %test.addr, align 8
     // SPIR64: store i32 8, ptr %s, align 4
     // AMDGCN: store ptr addrspace(1) %test, ptr addrspace(5) %test.addr, align 8
     // AMDGCN: store i32 8, ptr addrspace(5) %s, align 4
 }

llvm/docs/SPIRVUsage.rst

.. table:: SPIR-V Environments
      Environment           Description
      ===================== ==============================================================
      *<empty>*/``unknown``  Defaults to the OpenCL environment.
      ===================== ==============================================================
 
 Example:
 
 ``-target spirv64v1.0`` can be used to compile for SPIR-V version 1.0 with 64-bit pointer width.
+
+.. _spirv-types:
+
+Representing special types in SPIR-V
+====================================
+
+SPIR-V specifies several kinds of opaque types. These types are represented
+using target extension types. These types are represented as follows:
+
+  .. table:: SPIR-V Opaque Types
+
+     ================== ====================== =========================================================================================
+     SPIR-V Type        LLVM type name         LLVM type arguments
+     ================== ====================== =========================================================================================
+     OpTypeImage        ``spirv.Image``        sampled type, dimensionality, depth, arrayed, MS, sampled, image format, access qualifier
+     OpTypeSampler      ``spirv.Sampler``      (none)
+     OpTypeSampledImage ``spirv.SampledImage`` sampled type, dimensionality, depth, arrayed, MS, sampled, image format, access qualifier
+     OpTypeEvent        ``spirv.Event``        (none)
+     OpTypeDeviceEvent  ``spirv.DeviceEvent``  (none)
+     OpTypeReserveId    ``spirv.ReserveId``    (none)
+     OpTypeQueue        ``spirv.Queue``        (none)
+     OpTypePipe         ``spirv.Pipe``         access qualifier
+     OpTypePipeStorage  ``spirv.PipeStorage``  (none)
+     ================== ====================== =========================================================================================
+
+All integer arguments take the same value as they do in their `corresponding
+SPIR-V instruction <https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#_type_declaration_instructions>`_.
+For example, the OpenCL type ``image2d_depth_ro_t`` would be represented in
+SPIR-V IR as ``target("spirv.Image", void, 1, 1, 0, 0, 0, 0, 0)``, with its
+dimensionality parameter as ``1`` meaning 2D. Sampled image types include the
+parameters of its underlying image type, so that a sampled image for the
+previous type has the representation
+``target("spirv.SampledImage, void, 1, 1, 0, 0, 0, 0, 0)``.