[SPIR-V] Add SPIR-V builtin functions and types

[bader]

Summary

• Make the OpenCL Builtin lookup agnostic to programming models
  
  Rename the OpenCL Builtin lookup system to make it agnostic. The Builtin emitter is now outputting the builtin information into a class so that different programming models can be produced and used.

• Add SPIR-V variants of TypeSampledImage as clang builtin type.
  
  This patch adds SPIR-V sampled image types as derivative of the builtin OpenCL Image types. For each OpenCL image type, clang defines a Sampled<image type> variant and lowered as a "spirv.SampledImage.<image>" llvm opaque type.

• Enable SPIR-V builtin lookup
  
  Add flag -fdeclare-spirv-builtins to enable lookup of SPIR-V builtins.
  
  If -fdeclare-spirv-builtins is passed to clang, the compiler will try to lookup for the builtin described in SPIRVBuiltins.td for any match. If a match is found, overloads are build for the match.
  
  This will enforce a stable way to express SPIR-V builtins and make them closer to how the translator mangles them. This will help ensuring builtin for CUDA does not break easily. This will also support any changes suggested by the SPIRV-LLVM people on how to represent builtins.
  
  Define SPIRV_BUILTIN_DECLARATIONS when passing -fdeclare-spirv-builtins to clang.
  
  Added OpenCL SPIR-V extended set builtins bindings and part of the core SPIR-V (mostly missing Images and Pipes)

TODO: Known vendor extensions are not implemented yet.

Signed-off-by: Victor Lomuller <[e-mail]>

[Anastasia]

Can you elaborate the overall design flow please? Particularly, it is not clear why we need to duplicate OpenCL types and builtins at source level for SPIR-V... Is there any way to avoid this?

clang/include/clang-c/Index.h

enum CXTypeKind {
   CXType_OCLIntelSubgroupAVCSicResult = 171,
   CXType_OCLIntelSubgroupAVCImeResultSingleRefStreamout = 172,
   CXType_OCLIntelSubgroupAVCImeResultDualRefStreamout = 173,
   CXType_OCLIntelSubgroupAVCImeSingleRefStreamin = 174,
 
   CXType_OCLIntelSubgroupAVCImeDualRefStreamin = 175,
 
   CXType_ExtVector = 176,
-  CXType_Atomic = 177
+  CXType_Atomic = 177,
+
+  /* SPIRV builtin types. */
+  CXType_SampledOCLImage1dRO = 178,
+  CXType_SampledOCLImage1dArrayRO = 179,
+  CXType_SampledOCLImage1dBufferRO = 180,
+  CXType_SampledOCLImage2dRO = 181,
+  CXType_SampledOCLImage2dArrayRO = 182,
+  CXType_SampledOCLImage2dDepthRO = 183,
+  CXType_SampledOCLImage2dArrayDepthRO = 184,
+  CXType_SampledOCLImage2dMSAARO = 185,
+  CXType_SampledOCLImage2dArrayMSAARO = 186,
+  CXType_SampledOCLImage2dMSAADepthRO = 187,
+  CXType_SampledOCLImage2dArrayMSAADepthRO = 188,
+  CXType_SampledOCLImage3dRO = 189
 };
 
 /**
  * Describes the calling convention of a function type
  */
 enum CXCallingConv {
   CXCallingConv_Default = 0,
   CXCallingConv_C = 1,

clang/include/clang/AST/ASTContext.h

public:
   CanQualType DependentTy, OverloadTy, BoundMemberTy, UnknownAnyTy;
   CanQualType BuiltinFnTy;
   CanQualType PseudoObjectTy, ARCUnbridgedCastTy;
   CanQualType ObjCBuiltinIdTy, ObjCBuiltinClassTy, ObjCBuiltinSelTy;
   CanQualType ObjCBuiltinBoolTy;
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
   CanQualType SingletonId;
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  CanQualType Sampled##SingletonId;
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
   CanQualType OCLSamplerTy, OCLEventTy, OCLClkEventTy;
   CanQualType OCLQueueTy, OCLReserveIDTy;
   CanQualType IncompleteMatrixIdxTy;
   CanQualType OMPArraySectionTy, OMPArrayShapingTy, OMPIteratorTy;
 #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \
   CanQualType Id##Ty;
 #include "clang/Basic/OpenCLExtensionTypes.def"
 #define SVE_TYPE(Name, Id, SingletonId) \

clang/include/clang/AST/Type.h

public:
   bool isAtomicType() const;                    // C11 _Atomic()
   bool isUndeducedAutoType() const;             // C++11 auto or
                                                 // C++14 decltype(auto)
   bool isTypedefNameType() const;               // typedef or alias template
 
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
   bool is##Id##Type() const;
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  bool isSampled##Id##Type() const;
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
 
   bool isImageType() const;                     // Any OpenCL image type
+  bool isSampledImageType() const;              // Any SPIR-V Sampled image type
 
   bool isSamplerT() const;                      // OpenCL sampler_t
   bool isEventT() const;                        // OpenCL event_t
   bool isClkEventT() const;                     // OpenCL clk_event_t
   bool isQueueT() const;                        // OpenCL queue_t
   bool isReserveIDT() const;                    // OpenCL reserve_id_t
 
 #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \

 /// This class is used for builtin types like 'int'.  Builtin
 /// types are always canonical and have a literal name field.
 class BuiltinType : public Type {
 public:
   enum Kind {
 // OpenCL image types
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) Id,
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) Sampled##Id,
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
 // OpenCL extension types
 #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) Id,
 #include "clang/Basic/OpenCLExtensionTypes.def"
 // SVE Types
 #define SVE_TYPE(Name, Id, SingletonId) Id,
 #include "clang/Basic/AArch64SVEACLETypes.def"
 // PPC MMA Types
 #define PPC_VECTOR_TYPE(Name, Id, Size) Id,

 }
 
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
   inline bool Type::is##Id##Type() const { \
     return isSpecificBuiltinType(BuiltinType::Id); \
   }
 #include "clang/Basic/OpenCLImageTypes.def"
 
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  inline bool Type::isSampled##Id##Type() const {                              \
+    return isSpecificBuiltinType(BuiltinType::Sampled##Id);                    \
+  }
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
+
 inline bool Type::isSamplerT() const {
   return isSpecificBuiltinType(BuiltinType::OCLSampler);
 }
 
 inline bool Type::isEventT() const {
   return isSpecificBuiltinType(BuiltinType::OCLEvent);
 }
 
 inline bool Type::isClkEventT() const {
   return isSpecificBuiltinType(BuiltinType::OCLClkEvent);
 }
 
 inline bool Type::isQueueT() const {
   return isSpecificBuiltinType(BuiltinType::OCLQueue);
 }
 
 inline bool Type::isReserveIDT() const {
   return isSpecificBuiltinType(BuiltinType::OCLReserveID);
 }
 
 inline bool Type::isImageType() const {
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) is##Id##Type() ||
+  return isSampledImageType() ||
+#include "clang/Basic/OpenCLImageTypes.def"
+         false; // end boolean or operation
+}
+
+inline bool Type::isSampledImageType() const {
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  isSampled##Id##Type() ||
   return
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
 #include "clang/Basic/OpenCLImageTypes.def"
       false; // end boolean or operation
 }
 
 inline bool Type::isPipeType() const {
   return isa<PipeType>(CanonicalType);
 }
 

clang/include/clang/AST/TypeProperties.td

let Class = BuiltinType in {
   }
 
   def : Creator<[{
       switch (kind) {
 #define IMAGE_TYPE(IMGTYPE, ID, SINGLETON_ID, ACCESS, SUFFIX) \
       case BuiltinType::ID: return ctx.SINGLETON_ID;
 #include "clang/Basic/OpenCLImageTypes.def"
 
+#define IMAGE_TYPE(IMGTYPE, ID, SINGLETON_ID, ACCESS, SUFFIX) \
+      case BuiltinType::Sampled##ID: return ctx.Sampled##SINGLETON_ID;
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
+
 #define EXT_OPAQUE_TYPE(EXTTYPE, ID, EXT) \
       case BuiltinType::ID: return ctx.ID##Ty;
 #include "clang/Basic/OpenCLExtensionTypes.def"
 
 #define SVE_TYPE(NAME, ID, SINGLETON_ID) \
       case BuiltinType::ID: return ctx.SINGLETON_ID;
 #include "clang/Basic/AArch64SVEACLETypes.def"
 

clang/include/clang/Basic/LangOptions.def

 LANGOPT(GPUAllowDeviceInit, 1, 0, "allowing device side global init functions for HIP")
 LANGOPT(GPUMaxThreadsPerBlock, 32, 1024, "default max threads per block for kernel launch bounds for HIP")
 LANGOPT(GPUDeferDiag, 1, 0, "defer host/device related diagnostic messages for CUDA/HIP")
 LANGOPT(GPUExcludeWrongSideOverloads, 1, 0, "always exclude wrong side overloads in overloading resolution for CUDA/HIP")
 
 LANGOPT(SYCLIsDevice      , 1, 0, "Generate code for SYCL device")
 LANGOPT(SYCLIsHost        , 1, 0, "SYCL host compilation")
 ENUM_LANGOPT(SYCLVersion  , SYCLMajorVersion, 2, SYCL_None, "Version of the SYCL standard used")
+LANGOPT(DeclareSPIRVBuiltins, 1, 0, "Declare SPIR-V builtin functions")
 
 LANGOPT(HIPUseNewLaunchAPI, 1, 0, "Use new kernel launching API for HIP")
 
 LANGOPT(SizedDeallocation , 1, 0, "sized deallocation")
 LANGOPT(AlignedAllocation , 1, 0, "aligned allocation")
 LANGOPT(AlignedAllocationUnavailable, 1, 0, "aligned allocation functions are unavailable")
 LANGOPT(NewAlignOverride  , 32, 0, "maximum alignment guaranteed by '::operator new(size_t)'")
 LANGOPT(ConceptSatisfactionCaching , 1, 1, "enable satisfaction caching for C++20 Concepts")

clang/include/clang/Driver/Options.td

def fdefault_calling_conv_EQ : Joined<["-"], "fdefault-calling-conv=">,
   NormalizedValues<["DCC_CDecl", "DCC_FastCall", "DCC_StdCall", "DCC_VectorCall", "DCC_RegCall"]>,
   MarshallingInfoEnum<LangOpts<"DefaultCallingConv">, "DCC_None">;
 
 // These options cannot be marshalled, because they are used to set up the LangOptions defaults.
 def finclude_default_header : Flag<["-"], "finclude-default-header">,
   HelpText<"Include default header file for OpenCL">;
 def fdeclare_opencl_builtins : Flag<["-"], "fdeclare-opencl-builtins">,
   HelpText<"Add OpenCL builtin function declarations (experimental)">;
+def fdeclare_spirv_builtins : Flag<["-"], "fdeclare-spirv-builtins">,
+  HelpText<"Add SPIR-V builtin function declarations (experimental)">;
 
 def fpreserve_vec3_type : Flag<["-"], "fpreserve-vec3-type">,
   HelpText<"Preserve 3-component vector type">,
   MarshallingInfoFlag<CodeGenOpts<"PreserveVec3Type">>;
 def fwchar_type_EQ : Joined<["-"], "fwchar-type=">,
   HelpText<"Select underlying type for wchar_t">, Values<"char,short,int">,
   NormalizedValues<["1", "2", "4"]>,
   MarshallingInfoEnum<LangOpts<"WCharSize">, "0">;

clang/include/clang/Serialization/ASTBitCodes.h

enum PredefinedTypeIDs {
 
   /// \brief The '__bf16' type
   PREDEF_TYPE_BFLOAT16_ID = 73,
 
 /// OpenCL image types with auto numeration
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
   PREDEF_TYPE_##Id##_ID,
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  PREDEF_TYPE_SAMPLED_##Id##_ID,
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
 /// \brief OpenCL extension types with auto numeration
 #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) PREDEF_TYPE_##Id##_ID,
 #include "clang/Basic/OpenCLExtensionTypes.def"
 // \brief SVE types with auto numeration
 #define SVE_TYPE(Name, Id, SingletonId) PREDEF_TYPE_##Id##_ID,
 #include "clang/Basic/AArch64SVEACLETypes.def"
 // \brief  PowerPC MMA types with auto numeration
 #define PPC_VECTOR_TYPE(Name, Id, Size) PREDEF_TYPE_##Id##_ID,

clang/lib/AST/ASTContext.cpp

void ASTContext::InitBuiltinTypes(const TargetInfo &Target,
   LongDoubleComplexTy = getComplexType(LongDoubleTy);
   Float128ComplexTy   = getComplexType(Float128Ty);
 
   // Builtin types for 'id', 'Class', and 'SEL'.
   InitBuiltinType(ObjCBuiltinIdTy, BuiltinType::ObjCId);
   InitBuiltinType(ObjCBuiltinClassTy, BuiltinType::ObjCClass);
   InitBuiltinType(ObjCBuiltinSelTy, BuiltinType::ObjCSel);
 
-  if (LangOpts.OpenCL) {
+  if (LangOpts.OpenCL || LangOpts.SYCLIsDevice) {
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
     InitBuiltinType(SingletonId, BuiltinType::Id);
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  InitBuiltinType(Sampled##SingletonId, BuiltinType::Sampled##Id);
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
 
     InitBuiltinType(OCLSamplerTy, BuiltinType::OCLSampler);
     InitBuiltinType(OCLEventTy, BuiltinType::OCLEvent);
     InitBuiltinType(OCLClkEventTy, BuiltinType::OCLClkEvent);
     InitBuiltinType(OCLQueueTy, BuiltinType::OCLQueue);
     InitBuiltinType(OCLReserveIDTy, BuiltinType::OCLReserveID);
 
 #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \
case Type::Builtin:
     case BuiltinType::OCLSampler:
     case BuiltinType::OCLEvent:
     case BuiltinType::OCLClkEvent:
     case BuiltinType::OCLQueue:
     case BuiltinType::OCLReserveID:
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
     case BuiltinType::Id:
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  case BuiltinType::Sampled##Id:
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
 #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \
   case BuiltinType::Id:
 #include "clang/Basic/OpenCLExtensionTypes.def"
       AS = getTargetAddressSpace(
           Target->getOpenCLTypeAddrSpace(getOpenCLTypeKind(T)));
       Width = Target->getPointerWidth(AS);
       Align = Target->getPointerAlign(AS);
       break;
if (!BT) {
     return OCLTK_Default;
   }
 
   switch (BT->getKind()) {
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
   case BuiltinType::Id:                                                        \
     return OCLTK_Image;
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  case BuiltinType::Sampled##Id:                                               \
+    return OCLTK_Image;
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
 
   case BuiltinType::OCLClkEvent:
     return OCLTK_ClkEvent;
 
   case BuiltinType::OCLEvent:
     return OCLTK_Event;
 
   case BuiltinType::OCLQueue:
#include "clang/Basic/RISCVVTypes.def"
     case BuiltinType::ObjCClass:
     case BuiltinType::ObjCSel:
       llvm_unreachable("@encoding ObjC primitive type");
 
     // OpenCL and placeholder types don't need @encodings.
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
     case BuiltinType::Id:
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  case BuiltinType::Sampled##Id:
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
 #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \
     case BuiltinType::Id:
 #include "clang/Basic/OpenCLExtensionTypes.def"
     case BuiltinType::OCLEvent:
     case BuiltinType::OCLClkEvent:
     case BuiltinType::OCLQueue:
     case BuiltinType::OCLReserveID:
     case BuiltinType::OCLSampler:

clang/lib/AST/ASTImporter.cpp

 }
 
 ExpectedType ASTNodeImporter::VisitBuiltinType(const BuiltinType *T) {
   switch (T->getKind()) {
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
   case BuiltinType::Id: \
     return Importer.getToContext().SingletonId;
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  case BuiltinType::Sampled##Id:                                               \
+    return Importer.getToContext().Sampled##SingletonId;
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
 #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \
   case BuiltinType::Id: \
     return Importer.getToContext().Id##Ty;
 #include "clang/Basic/OpenCLExtensionTypes.def"
 #define SVE_TYPE(Name, Id, SingletonId) \
   case BuiltinType::Id: \
     return Importer.getToContext().SingletonId;
 #include "clang/Basic/AArch64SVEACLETypes.def"

clang/lib/AST/ExprConstant.cpp

#include "clang/AST/BuiltinTypes.def"
     case BuiltinType::NullPtr:
 
     case BuiltinType::ObjCId:
     case BuiltinType::ObjCClass:
     case BuiltinType::ObjCSel:
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
     case BuiltinType::Id:
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  case BuiltinType::Sampled##Id:
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
 #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \
     case BuiltinType::Id:
 #include "clang/Basic/OpenCLExtensionTypes.def"
     case BuiltinType::OCLSampler:
     case BuiltinType::OCLEvent:
     case BuiltinType::OCLClkEvent:
     case BuiltinType::OCLQueue:
     case BuiltinType::OCLReserveID:

clang/lib/AST/ItaniumMangle.cpp

case BuiltinType::ObjCSel:
     Out << "13objc_selector";
     break;
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
   case BuiltinType::Id: \
     type_name = "ocl_" #ImgType "_" #Suffix; \
     Out << type_name.size() << type_name; \
     break;
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  case BuiltinType::Sampled##Id:                                               \
+    type_name = "__spirv_SampledImage__" #ImgType "_" #Suffix;                 \
+    Out << type_name.size() << type_name;                                      \
+    break;
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
   case BuiltinType::OCLSampler:
     Out << "11ocl_sampler";
     break;
   case BuiltinType::OCLEvent:
     Out << "9ocl_event";
     break;
   case BuiltinType::OCLClkEvent:
     Out << "12ocl_clkevent";

clang/lib/AST/MicrosoftMangle.cpp

case BuiltinType::ObjCSel:
     mangleArtificialTagType(TTK_Struct, "objc_selector");
     break;
 
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
   case BuiltinType::Id: \
     Out << "PAUocl_" #ImgType "_" #Suffix "@@"; \
     break;
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  case BuiltinType::Sampled##Id:                                               \
+    Out << "PAU__spirv_SampledImage__" #ImgType "_" #Suffix "@@";              \
+    break;
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
   case BuiltinType::OCLSampler:
     Out << "PA";
     mangleArtificialTagType(TTK_Struct, "ocl_sampler");
     break;
   case BuiltinType::OCLEvent:
     Out << "PA";
     mangleArtificialTagType(TTK_Struct, "ocl_event");
     break;

clang/lib/AST/NSAPI.cpp

NSAPI::getNSNumberFactoryMethodKind(QualType T) const {
   case BuiltinType::Float128:
   case BuiltinType::NullPtr:
   case BuiltinType::ObjCClass:
   case BuiltinType::ObjCId:
   case BuiltinType::ObjCSel:
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
   case BuiltinType::Id:
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  case BuiltinType::Sampled##Id:
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
 #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \
   case BuiltinType::Id:
 #include "clang/Basic/OpenCLExtensionTypes.def"
   case BuiltinType::OCLSampler:
   case BuiltinType::OCLEvent:
   case BuiltinType::OCLClkEvent:
   case BuiltinType::OCLQueue:
   case BuiltinType::OCLReserveID:

clang/lib/AST/PrintfFormatString.cpp

bool PrintfSpecifier::fixType(QualType QT, const LangOptions &LangOpt,
   case BuiltinType::SatUFract:
   case BuiltinType::SatULongFract:
     // Various types which are non-trivial to correct.
     return false;
 
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
   case BuiltinType::Id:
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  case BuiltinType::Sampled##Id:
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
 #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \
   case BuiltinType::Id:
 #include "clang/Basic/OpenCLExtensionTypes.def"
 #define SVE_TYPE(Name, Id, SingletonId) \
   case BuiltinType::Id:
 #include "clang/Basic/AArch64SVEACLETypes.def"
 #define PPC_VECTOR_TYPE(Name, Id, Size) \
   case BuiltinType::Id:

clang/lib/AST/Type.cpp

StringRef BuiltinType::getName(const PrintingPolicy &Policy) const {
   case ObjCClass:
     return "Class";
   case ObjCSel:
     return "SEL";
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
   case Id: \
     return "__" #Access " " #ImgType "_t";
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  case Sampled##Id:                                                            \
+    return "__ocl_sampled_" #ImgType "_" #Suffix "_t";
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
   case OCLSampler:
     return "sampler_t";
   case OCLEvent:
     return "event_t";
   case OCLClkEvent:
     return "clk_event_t";
   case OCLQueue:
     return "queue_t";
#include "clang/AST/BuiltinTypes.def"
 
     case BuiltinType::Void:
     case BuiltinType::ObjCId:
     case BuiltinType::ObjCClass:
     case BuiltinType::ObjCSel:
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
     case BuiltinType::Id:
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  case BuiltinType::Sampled##Id:
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
 #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \
     case BuiltinType::Id:
 #include "clang/Basic/OpenCLExtensionTypes.def"
     case BuiltinType::OCLSampler:
     case BuiltinType::OCLEvent:
     case BuiltinType::OCLClkEvent:
     case BuiltinType::OCLQueue:
     case BuiltinType::OCLReserveID:

clang/lib/AST/TypeLoc.cpp

TypeSpecifierType BuiltinTypeLoc::getWrittenTypeSpec() const {
   case BuiltinType::ARCUnbridgedCast:
   case BuiltinType::PseudoObject:
   case BuiltinType::ObjCId:
   case BuiltinType::ObjCClass:
   case BuiltinType::ObjCSel:
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
   case BuiltinType::Id:
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  case BuiltinType::Sampled##Id:
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
 #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \
   case BuiltinType::Id:
 #include "clang/Basic/OpenCLExtensionTypes.def"
   case BuiltinType::OCLSampler:
   case BuiltinType::OCLEvent:
   case BuiltinType::OCLClkEvent:
   case BuiltinType::OCLQueue:
   case BuiltinType::OCLReserveID:

clang/lib/CodeGen/CGDebugInfo.h

class CGDebugInfo {
   llvm::MDNode *CurInlinedAt = nullptr;
   llvm::DIType *VTablePtrType = nullptr;
   llvm::DIType *ClassTy = nullptr;
   llvm::DICompositeType *ObjTy = nullptr;
   llvm::DIType *SelTy = nullptr;
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
   llvm::DIType *SingletonId = nullptr;
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  llvm::DIType *Sampled##SingletonId = nullptr;
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
   llvm::DIType *OCLSamplerDITy = nullptr;
   llvm::DIType *OCLEventDITy = nullptr;
   llvm::DIType *OCLClkEventDITy = nullptr;
   llvm::DIType *OCLQueueDITy = nullptr;
   llvm::DIType *OCLNDRangeDITy = nullptr;
   llvm::DIType *OCLReserveIDDITy = nullptr;
 #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \
   llvm::DIType *Id##Ty = nullptr;

clang/lib/CodeGen/CGDebugInfo.cpp

case BuiltinType::ObjCSel: {
     return SelTy;
   }
 
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
   case BuiltinType::Id:                                                        \
     return getOrCreateStructPtrType("opencl_" #ImgType "_" #Suffix "_t",       \
                                     SingletonId);
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  case BuiltinType::Sampled##Id:                                               \
+    return getOrCreateStructPtrType(                                           \
+        "spirv_sampled_" #ImgType "_" #Suffix "_t", Sampled##SingletonId);
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
   case BuiltinType::OCLSampler:
     return getOrCreateStructPtrType("opencl_sampler_t", OCLSamplerDITy);
   case BuiltinType::OCLEvent:
     return getOrCreateStructPtrType("opencl_event_t", OCLEventDITy);
   case BuiltinType::OCLClkEvent:
     return getOrCreateStructPtrType("opencl_clk_event_t", OCLClkEventDITy);
   case BuiltinType::OCLQueue:
     return getOrCreateStructPtrType("opencl_queue_t", OCLQueueDITy);

clang/lib/CodeGen/CGOpenCLRuntime.cpp

default:
     llvm_unreachable("Unexpected opencl builtin type!");
     return nullptr;
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
   case BuiltinType::Id: \
     return llvm::PointerType::get( \
         llvm::StructType::create(Ctx, "opencl." #ImgType "_" #Suffix "_t"), \
         AddrSpc);
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  case BuiltinType::Sampled##Id:                                               \
+    return llvm::PointerType::get(                                             \
+        llvm::StructType::create(Ctx, "spirv.SampledImage." #ImgType           \
+                                      "_" #Suffix "_t"),                       \
+        AddrSpc);
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
   case BuiltinType::OCLSampler:
     return getSamplerType(T);
   case BuiltinType::OCLEvent:
     return llvm::PointerType::get(
         llvm::StructType::create(Ctx, "opencl.event_t"), AddrSpc);
   case BuiltinType::OCLClkEvent:
     return llvm::PointerType::get(
         llvm::StructType::create(Ctx, "opencl.clk_event_t"), AddrSpc);

clang/lib/CodeGen/CodeGenTypes.cpp

case Type::Builtin: {
     case BuiltinType::UInt128:
     case BuiltinType::Int128:
       ResultType = llvm::IntegerType::get(getLLVMContext(), 128);
       break;
 
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
     case BuiltinType::Id:
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  case BuiltinType::Sampled##Id:
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
 #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \
     case BuiltinType::Id:
 #include "clang/Basic/OpenCLExtensionTypes.def"
     case BuiltinType::OCLSampler:
     case BuiltinType::OCLEvent:
     case BuiltinType::OCLClkEvent:
     case BuiltinType::OCLQueue:
     case BuiltinType::OCLReserveID:

clang/lib/CodeGen/ItaniumCXXABI.cpp

switch (Ty->getKind()) {
     case BuiltinType::Char32:
     case BuiltinType::Int128:
     case BuiltinType::UInt128:
       return true;
 
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
     case BuiltinType::Id:
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  case BuiltinType::Sampled##Id:
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
 #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \
     case BuiltinType::Id:
 #include "clang/Basic/OpenCLExtensionTypes.def"
     case BuiltinType::OCLSampler:
     case BuiltinType::OCLEvent:
     case BuiltinType::OCLClkEvent:
     case BuiltinType::OCLQueue:
     case BuiltinType::OCLReserveID:

clang/lib/Frontend/CompilerInvocation.cpp

if (const Arg *A = Args.getLastArg(OPT_cl_std_EQ)) {
     else
       LangStd = OpenCLLangStd;
   }
 
   // These need to be parsed now. They are used to set OpenCL defaults.
   Opts.IncludeDefaultHeader = Args.hasArg(OPT_finclude_default_header);
   Opts.DeclareOpenCLBuiltins = Args.hasArg(OPT_fdeclare_opencl_builtins);
 
+  Opts.SYCLIsDevice = Args.hasArg(options::OPT_fsycl_is_device);
+  Opts.DeclareSPIRVBuiltins = Args.hasArg(OPT_fdeclare_spirv_builtins);
+
   CompilerInvocation::setLangDefaults(Opts, IK, T, Includes, LangStd);
 
   // The key paths of codegen options defined in Options.td start with
   // "LangOpts->". Let's provide the expected variable name and type.
   LangOptions *LangOpts = &Opts;
 
 #define LANG_OPTION_WITH_MARSHALLING(                                          \
     PREFIX_TYPE, NAME, ID, KIND, GROUP, ALIAS, ALIASARGS, FLAGS, PARAM,        \

clang/lib/Frontend/InitPreprocessor.cpp

static void InitializeStandardPredefinedMacros(const TargetInfo &TI,
   if (LangOpts.SYCLIsDevice || LangOpts.SYCLIsHost) {
     // SYCL Version is set to a value when building SYCL applications
     if (LangOpts.getSYCLVersion() == LangOptions::SYCL_2017)
       Builder.defineMacro("CL_SYCL_LANGUAGE_VERSION", "121");
     else if (LangOpts.getSYCLVersion() == LangOptions::SYCL_2020)
       Builder.defineMacro("SYCL_LANGUAGE_VERSION", "202001");
   }
 
+  if (LangOpts.DeclareSPIRVBuiltins) {
+    Builder.defineMacro("__SPIRV_BUILTIN_DECLARATIONS__");
+  }
+
   // Not "standard" per se, but available even with the -undef flag.
   if (LangOpts.AsmPreprocessor)
     Builder.defineMacro("__ASSEMBLER__");
   if (LangOpts.CUDA && !LangOpts.HIP)
     Builder.defineMacro("__CUDA__");
   if (LangOpts.HIP) {
     Builder.defineMacro("__HIP__");
     Builder.defineMacro("__HIPCC__");

clang/lib/Index/USRGeneration.cpp

if (const BuiltinType *BT = T->getAs<BuiltinType>()) {
           c = 'n'; break;
 #define BUILTIN_TYPE(Id, SingletonId)
 #define PLACEHOLDER_TYPE(Id, SingletonId) case BuiltinType::Id:
 #include "clang/AST/BuiltinTypes.def"
         case BuiltinType::Dependent:
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
         case BuiltinType::Id:
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  case BuiltinType::Sampled##Id:
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
 #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \
         case BuiltinType::Id:
 #include "clang/Basic/OpenCLExtensionTypes.def"
         case BuiltinType::OCLEvent:
         case BuiltinType::OCLClkEvent:
         case BuiltinType::OCLQueue:
         case BuiltinType::OCLReserveID:
         case BuiltinType::OCLSampler:

clang/lib/Sema/CMakeLists.txt

 set(LLVM_LINK_COMPONENTS
   Core
   FrontendOpenMP
   MC
   Support
   )
 
 clang_tablegen(OpenCLBuiltins.inc -gen-clang-opencl-builtins
   SOURCE OpenCLBuiltins.td
   TARGET ClangOpenCLBuiltinsImpl
   )
 
+clang_tablegen(SPIRVBuiltins.inc -gen-clang-spirv-builtins
+ SOURCE SPIRVBuiltins.td
+ TARGET ClangSPIRVBuiltinsImpl
+ )
+
 add_clang_library(clangSema
   AnalysisBasedWarnings.cpp
   CodeCompleteConsumer.cpp
   DeclSpec.cpp
   DelayedDiagnostic.cpp
   IdentifierResolver.cpp
   JumpDiagnostics.cpp
   MultiplexExternalSemaSource.cpp
add_clang_library(clangSema
   SemaTemplateInstantiate.cpp
   SemaTemplateInstantiateDecl.cpp
   SemaTemplateVariadic.cpp
   SemaType.cpp
   TypeLocBuilder.cpp
 
   DEPENDS
   ClangOpenCLBuiltinsImpl
+  ClangSPIRVBuiltinsImpl
   omp_gen
 
   LINK_LIBS
   clangAST
   clangAnalysis
   clangBasic
   clangEdit
   clangLex
   )

clang/lib/Sema/OpenCLBuiltins.td

class Builtin<string _Name, list<Type> _Signature, list<bit> _Attributes = Attr.None> {
   // (the return type).
   list<Type> Signature = _Signature;
   // Function attribute __attribute__((pure))
   bit IsPure = _Attributes[0];
   // Function attribute __attribute__((const))
   bit IsConst = _Attributes[1];
   // Function attribute __attribute__((convergent))
   bit IsConv = _Attributes[2];
+  // Is function a variadic one
+  bit IsVariadic = 0;
   // OpenCL extensions to which the function belongs.
   FunctionExtension Extension = FuncExtNone;
   // Version of OpenCL from which the function is available (e.g.: CL10).
   // MinVersion is inclusive.
   Version MinVersion = CL10;
   // Version of OpenCL from which the function is not supported anymore.
   // MaxVersion is exclusive.
   // CLAll makes the function available for all versions.

clang/lib/Sema/SPIRVBuiltins.td

+//==--- SPIRVBuiltins.td - SPIRV builtin declarations -------------------===//
+//
+//                     The LLVM Compiler Infrastructure
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file contains TableGen definitions for SPIR-V builtin function
+// declarations.  In case of an unresolved function name, Clang will check for
+// a function described in this file when -fdeclare-spirv-builtins is specified.
+//
+//===----------------------------------------------------------------------===//
+
+//===----------------------------------------------------------------------===//
+//              Definitions of miscellaneous basic entities.
+//===----------------------------------------------------------------------===//
+// TODO: basic entities declaration with OpenCLBuiltins.td
+
+// TODO: Manage version using the JSON grammar.  Unused for now.
+class Version<int _Version> {
+  int ID = _Version;
+}
+def SPIRVAll : Version<  0>;
+
+// Address spaces
+// Pointer types need to be assigned an address space.
+class AddressSpace<string _AS> {
+  string Name = _AS;
+}
+// Default is important for the frontend as there is not necessarily
+// an automatic conversion from this address space to
+// the one it will be lowered to.
+// This file assumes it will get lowered to generic or private.
+def DefaultAS    : AddressSpace<"clang::LangAS::Default">;
+def PrivateAS    : AddressSpace<"clang::LangAS::sycl_private">;
+def GlobalAS     : AddressSpace<"clang::LangAS::sycl_global">;
+def ConstantAS   : AddressSpace<"clang::LangAS::opencl_constant">;
+def LocalAS      : AddressSpace<"clang::LangAS::sycl_local">;
+def GenericAS    : AddressSpace<"clang::LangAS::opencl_generic">;
+
+// TODO: Manage capabilities. Unused for now.
+class AbstractExtension<string _Ext> {
+  string ExtName = _Ext;
+}
+
+// Extension associated to a builtin function.
+class FunctionExtension<string _Ext> : AbstractExtension<_Ext>;
+
+// FunctionExtension definitions.
+def FuncExtNone                          : FunctionExtension<"">;
+
+// Extension associated to a type.  This enables implicit conditionalization of
+// builtin function overloads containing a type that depends on an extension.
+// During overload resolution, when a builtin function overload contains a type
+// with a TypeExtension, those overloads are skipped when the extension is
+// disabled.
+class TypeExtension<string _Ext> : AbstractExtension<_Ext>;
+
+// TypeExtension definitions.
+def NoTypeExt   : TypeExtension<"">;
+
+// Qualified Type.  These map to ASTContext::QualType.
+// TODO: Create a QualTypeFromASTContext.
+//       To fully make sense here, this class should represent
+//       the QualType only. How the QualType is accessed should be separated.
+class QualType<string _TypeExpr, bit _IsAbstract=0, bit _IsSigned=0> {
+  // Name of the field or function in a clang::ASTContext
+  // E.g. Name="IntTy" for the int type, and "getIntPtrType()" for an intptr_t
+  string TypeExpr = _TypeExpr;
+  // Some QualTypes in this file represent an abstract type for which there is
+  // no corresponding AST QualType, e.g. a GenType or an `image2d_t` type
+  // without access qualifiers.
+  bit IsAbstract = _IsAbstract;
+  bit IsSigned = _IsSigned;
+}
+
+// Qualified Type.  These map to a function taking an ASTContext
+// and returning a QualType.
+// Instead of direclty accessing ASTContext fields, the builtin lookup can
+// call a function to extract the correct type for the call.
+// The name will be interpreted as the function to call
+// rather than the field to access.
+class QualTypeFromFunction<string _Name, bit _IsAbstract=0, bit _IsSigned=0> :
+      QualType<_Name, _IsAbstract, _IsSigned> {
+// TODO: At the moment the user is expected to write the function outside this file.
+//       Although they could be generated in the .inc file and
+//       the user would only have to provide the body here
+//       (like it can be done for attributes for instance).
+}
+
+// List of integers.
+class IntList<string _Name, list<int> _List> {
+  string Name = _Name;
+  list<int> List = _List;
+}
+
+// Basic data types (int, float, image2d_t, ...).
+// Its child classes can represent concrete types (e.g. VectorType) or
+// abstract types (e.g. GenType).
+class Type<string _Name, QualType _QTExpr> {
+  // Name of the Type.
+  string Name = _Name;
+  // QualType associated with this type.
+  QualType QTExpr = _QTExpr;
+  // Size of the vector (if applicable).
+  int VecWidth = 1;
+  // Size of the element in bits.
+  int ElementSize = 1;
+  // Is a integer.
+  bit IsInteger = 0;
+  // Is a signed integer.
+  bit IsSigned = 1;
+  // Is a float.
+  bit IsFloat = 0;
+  // Is a pointer.
+  bit IsPointer = 0;
+  // "const" qualifier.
+  bit IsConst = 0;
+  // "volatile" qualifier.
+  bit IsVolatile = 0;
+  // Access qualifier. Must be one of ("RO", "WO", "RW").
+  string AccessQualifier = "";
+  // Address space.
+  string AddrSpace = DefaultAS.Name;
+  // Extension that needs to be enabled to expose a builtin that uses this type.
+  TypeExtension Extension = NoTypeExt;
+}
+
+class FundamentalType<string _Name, QualType _QTName, int _Size> : Type<_Name, _QTName> {
+  // Inherited fields
+  let ElementSize = _Size;
+}
+
+// Integer Type.
+class IntType<string _Name, QualType _QTName, int _Size> : FundamentalType<_Name, _QTName, _Size> {
+  // Inherited fields
+  let IsInteger = 1;
+  let IsSigned = 1;
+}
+
+// Unsigned integer Type.
+class UIntType<string _Name, QualType _QTName, int _Size> : FundamentalType<_Name, _QTName, _Size> {
+  // Inherited fields
+  let IsInteger = 1;
+  let IsSigned = 0;
+}
+
+// Floating Type.
+class FPType<string _Name, QualType _QTName, int _Size> : FundamentalType<_Name, _QTName, _Size> {
+  // Inherited fields
+  let IsFloat = 1;
+}
+
+class CompoundType<Type _Ty> : Type<_Ty.Name, _Ty.QTExpr> {
+  // Inherited fields
+  let VecWidth = _Ty.VecWidth;
+  let ElementSize = _Ty.ElementSize;
+  let IsInteger = _Ty.IsInteger;
+  let IsSigned = _Ty.IsSigned;
+  let IsFloat = _Ty.IsFloat;
+  let IsPointer = _Ty.IsPointer;
+  let IsConst = _Ty.IsConst;
+  let IsVolatile = _Ty.IsVolatile;
+  let AccessQualifier = _Ty.AccessQualifier;
+  let AddrSpace = _Ty.AddrSpace;
+  let Extension = _Ty.Extension;
+
+  Type ElementType = _Ty;
+}
+
+// Vector types (e.g. int2, int3, int16, float8, ...).
+class VectorType<Type _Ty, int _VecWidth> : Type<_Ty.Name, _Ty.QTExpr> {
+  let VecWidth = _VecWidth;
+  let AccessQualifier = "";
+  // Inherited fields
+  let ElementSize = _Ty.ElementSize;
+  let IsInteger = _Ty.IsInteger;
+  let IsSigned = _Ty.IsSigned;
+  let IsFloat = _Ty.IsFloat;
+  let IsPointer = _Ty.IsPointer;
+  let IsConst = _Ty.IsConst;
+  let IsVolatile = _Ty.IsVolatile;
+  let AccessQualifier = _Ty.AccessQualifier;
+  let AddrSpace = _Ty.AddrSpace;
+  let Extension = _Ty.Extension;
+}
+
+// Pointer types (e.g. int*, float*, ...).
+class PointerType<Type _Ty, AddressSpace _AS = DefaultAS> :
+    CompoundType<_Ty> {
+  // Inherited fields
+  let IsPointer = 1;
+  let AddrSpace = _AS.Name;
+  let Extension = _Ty.Extension;
+}
+
+// Const types (e.g. const int).
+class ConstType<Type _Ty> : CompoundType<_Ty> {
+  // Inherited fields
+  let IsConst = 1;
+  let Extension = _Ty.Extension;
+}
+
+// Volatile types (e.g. volatile int).
+class VolatileType<Type _Ty> : CompoundType<_Ty> {
+  // Inherited fields
+  let IsVolatile = 1;
+  let Extension = _Ty.Extension;
+}
+
+// Image types (e.g. image2d).
+class ImageType<Type _Ty, string _AccessQualifier> :
+    Type<_Ty.Name, QualType<_Ty.QTExpr.TypeExpr # _AccessQualifier # "Ty", 0>> {
+  let VecWidth = 0;
+  let AccessQualifier = _AccessQualifier;
+  // Inherited fields
+  let ElementSize = _Ty.ElementSize;
+  let IsInteger = _Ty.IsInteger;
+  let IsSigned = _Ty.IsSigned;
+  let IsFloat = _Ty.IsFloat;
+  let IsPointer = _Ty.IsPointer;
+  let IsConst = _Ty.IsConst;
+  let IsVolatile = _Ty.IsVolatile;
+  let AddrSpace = _Ty.AddrSpace;
+  let Extension = _Ty.Extension;
+}
+
+// List of Types.
+class TypeList<list<Type> _Type> {
+  list<Type> List = _Type;
+}
+
+// A GenericType is an abstract type that defines a set of types as a
+// combination of Types and vector sizes.
+//
+// For example, if TypeList = <int, float> and VectorList = <1, 2, 4>, then it
+// represents <int, int2, int4, float, float2, float4>.
+//
+// Some rules apply when using multiple GenericType arguments in a declaration:
+//   1. The number of vector sizes must be equal or 1 for all gentypes in a
+//      declaration.
+//   2. The number of Types must be equal or 1 for all gentypes in a
+//      declaration.
+//   3. Generic types are combined by iterating over all generic types at once.
+//      For example, for the following GenericTypes
+//        GenT1 = GenericType<half, [1, 2]> and
+//        GenT2 = GenericType<float, int, [1, 2]>
+//      A declaration f(GenT1, GenT2) results in the combinations
+//        f(half, float), f(half2, float2), f(half, int), f(half2, int2) .
+//   4. "sgentype" from the OpenCL specification is supported by specifying
+//      a single vector size.
+//      For example, for the following GenericTypes
+//        GenT = GenericType<half, int, [1, 2]> and
+//        SGenT = GenericType<half, int, [1]>
+//      A declaration f(GenT, SGenT) results in the combinations
+//        f(half, half), f(half2, half), f(int, int), f(int2, int) .
+class GenericType<string _Ty, TypeList _TypeList, IntList _VectorList> :
+    Type<_Ty, QualType<"null", 1>> {
+  // Possible element types of the generic type.
+  TypeList TypeList = _TypeList;
+  // Possible vector sizes of the types in the TypeList.
+  IntList VectorList = _VectorList;
+  // The VecWidth field is ignored for GenericTypes. Use VectorList instead.
+  let VecWidth = 0;
+}
+
+// Builtin function attributes.
+def Attr {
+  list<bit> None = [0, 0, 0];
+  list<bit> Pure = [1, 0, 0];
+  list<bit> Const = [0, 1, 0];
+  list<bit> Convergent = [0, 0, 1];
+}
+
+//===----------------------------------------------------------------------===//
+//                      Class for builtin functions
+//===----------------------------------------------------------------------===//
+class Builtin<string _Name, list<Type> _Signature, list<bit> _Attributes = Attr.None> {
+  // Name of the builtin function
+  string Name = _Name;
+  // List of types used by the function. The first one is the return type and
+  // the following are the arguments. The list must have at least one element
+  // (the return type).
+  list<Type> Signature = _Signature;
+  // Function attribute __attribute__((pure))
+  bit IsPure = _Attributes[0];
+  // Function attribute __attribute__((const))
+  bit IsConst = _Attributes[1];
+  // Function attribute __attribute__((convergent))
+  bit IsConv = _Attributes[2];
+  // Is function a variadic one
+  bit IsVariadic = 0;
+  // OpenCL extensions to which the function belongs.
+  FunctionExtension Extension = FuncExtNone;
+  // Version from which the function is available.
+  // MinVersion is inclusive.
+  Version MinVersion = SPIRVAll;
+  // Version from which the function is not supported anymore.
+  // MaxVersion is exclusive.
+  // SPIRVAll makes the function available for all versions.
+  Version MaxVersion = SPIRVAll;
+}
+
+// Helper to declare SPIR-V Core builtins.
+class SPVBuiltin<string _Name, list<Type> _Signature, list<bit> _Attributes = Attr.None> :
+  Builtin<"__spirv_" # _Name, _Signature, _Attributes> {}
+
+// Helper to declare OpenCL SPIR-V extended set builtins.
+class OCLSPVBuiltin<string _Name, list<Type> _Signature, list<bit> _Attributes = Attr.None> :
+  SPVBuiltin<"ocl_" # _Name, _Signature, _Attributes> {}
+
+class ConstOCLSPVBuiltin<string _Name, list<Type> _Signature> :
+  OCLSPVBuiltin<_Name, _Signature, Attr.Const> {}
+
+//===----------------------------------------------------------------------===//
+//                 Definitions of types
+//===----------------------------------------------------------------------===//
+
+// OpenCL v1.0/1.2/2.0 s6.1.1: Built-in Scalar Data Types.
+def Bool      : IntType<"bool",   QualType<"Context.BoolTy">, 1>;
+def Char      : IntType<"char",      QualType<"Context.CharTy", 0, 1>, 8>;
+def SChar     : IntType<"schar",      QualType<"Context.SignedCharTy", 0, 1>, 8>;
+def UChar     : UIntType<"uchar",     QualType<"Context.UnsignedCharTy">, 8>;
+def Short     : IntType<"short",     QualType<"Context.ShortTy", 0, 1>, 16>;
+def UShort    : UIntType<"ushort",    QualType<"Context.UnsignedShortTy">, 16>;
+def Int       : IntType<"int",       QualType<"Context.IntTy", 0, 1>, 32>;
+def UInt      : UIntType<"uint",      QualType<"Context.UnsignedIntTy">, 32>;
+def Long      : IntType<"long",      QualType<"Context.getIntTypeForBitwidth(64, true)", 0, 1>, 64>;
+def ULong     : UIntType<"ulong",     QualType<"Context.getIntTypeForBitwidth(64, false)">, 64>;
+def Float     : FPType<"float",     QualType<"Context.FloatTy">, 32>;
+def Double    : FPType<"double",    QualType<"Context.DoubleTy">, 64>;
+def Half      : FPType<"half",      QualTypeFromFunction<"GetFloat16Type">, 16>;
+def Void      : Type<"void",      QualType<"Context.VoidTy">>;
+// FIXME: ensure this is portable...
+def Size      : Type<"size_t",    QualType<"Context.getSizeType()">>;
+
+def Sampler               : Type<"sampler_t", QualType<"Context.OCLSamplerTy">>;
+def Event                 : Type<"event_t", QualType<"Context.OCLEventTy">>;
+
+//===----------------------------------------------------------------------===//
+//                 Definitions of gentype variants
+//===----------------------------------------------------------------------===//
+
+// Vector width lists.
+def VecAndScalar: IntList<"VecAndScalar", [1, 2, 3, 4, 8, 16]>;
+def VecNoScalar : IntList<"VecNoScalar", [2, 3, 4, 8, 16]>;
+def Vec1        : IntList<"Vec1", [1]>;
+def Vec2        : IntList<"Vec2", [2]>;
+def Vec4        : IntList<"Vec4", [4]>;
+def Vec8        : IntList<"Vec8", [8]>;
+def Vec16       : IntList<"Vec16", [16]>;
+def Vec1234     : IntList<"Vec1234", [1, 2, 3, 4]>;
+
+// Type lists.
+def TLAll           : TypeList<[Char, UChar, Short,  UShort, Int,  UInt, Long,  ULong, Float, Double, Half]>;
+def TLAllUnsigned   : TypeList<[UChar, UChar, UShort, UShort, UInt, UInt, ULong, ULong, UInt,  ULong,  UShort]>;
+def TLFloat         : TypeList<[Float, Double, Half]>;
+def TLSignedInts    : TypeList<[Char, Short, Int, Long]>;
+def TLUnsignedInts  : TypeList<[UChar, UShort, UInt, ULong]>;
+
+// Signed to Unsigned conversion
+def TLSToUSignedInts    : TypeList<[Char, Short, Int, Long]>;
+def TLSToUUnsignedInts  : TypeList<[UChar, UShort, UInt, ULong]>;
+
+def TLIntLongFloats : TypeList<[Int, UInt, Long, ULong, Float, Double, Half]>;
+
+// All unsigned integer types twice, to facilitate unsigned return types for e.g.
+// uchar abs(char) and
+// uchar abs(uchar).
+def TLAllUIntsTwice : TypeList<[UChar, UChar, UChar, UShort, UShort, UInt, UInt, ULong, ULong]>;
+
+def TLAllInts       : TypeList<[Char, SChar, UChar, Short, UShort, Int, UInt, Long, ULong]>;
+
+// GenType definitions for multiple base types (e.g. all floating point types,
+// or all integer types).
+// All types
+def AGenType1              : GenericType<"AGenType1", TLAll, Vec1>;
+def AGenTypeN              : GenericType<"AGenTypeN", TLAll, VecAndScalar>;
+def AGenTypeNNoScalar      : GenericType<"AGenTypeNNoScalar", TLAll, VecNoScalar>;
+// All integer
+def AIGenType1             : GenericType<"AIGenType1", TLAllInts, Vec1>;
+def AIGenTypeN             : GenericType<"AIGenTypeN", TLAllInts, VecAndScalar>;
+def AUIGenTypeN             : GenericType<"AUIGenTypeN", TLUnsignedInts, VecAndScalar>;
+def ASIGenTypeN             : GenericType<"ASIGenTypeN", TLSignedInts, VecAndScalar>;
+def AIGenTypeNNoScalar     : GenericType<"AIGenTypeNNoScalar", TLAllInts, VecNoScalar>;
+// All integer to unsigned
+def AI2UGenTypeN           : GenericType<"AI2UGenTypeN", TLAllUIntsTwice, VecAndScalar>;
+// Signed integer
+def SGenTypeN              : GenericType<"SGenTypeN", TLSignedInts, VecAndScalar>;
+// Unsigned integer
+def UGenTypeN              : GenericType<"UGenTypeN", TLUnsignedInts, VecAndScalar>;
+// Float
+def FGenTypeN              : GenericType<"FGenTypeN", TLFloat, VecAndScalar>;
+// (u)int, (u)long, and all floats
+def IntLongFloatGenType1   : GenericType<"IntLongFloatGenType1", TLIntLongFloats, Vec1>;
+
+// GenType definitions for every single base type (e.g. fp32 only).
+// Names are like: GenTypeFloatVecAndScalar.
+foreach Type = [Char, SChar, UChar, Short, UShort,
+                Int, UInt, Long, ULong,
+                Float, Double, Half] in {
+  foreach VecSizes = [VecAndScalar, VecNoScalar] in {
+    def "GenType" # Type # VecSizes :
+              GenericType<"GenType" # Type # VecSizes,
+                          TypeList<[Type]>, VecSizes>;
+  }
+}
+
+// GenType definitions for vec1234.
+foreach Type = [Float, Double, Half] in {
+  def "GenType" # Type # Vec1234 :
+              GenericType<"GenType" # Type # Vec1234,
+                          TypeList<[Type]>, Vec1234>;
+}
+
+//===----------------------------------------------------------------------===//
+//                 Definitions of builtins
+//           extinst.opencl.std.100.grammar.json
+//===----------------------------------------------------------------------===//
+
+// 2.1. Math extended instructions
+
+
+foreach name  = ["acos", "acosh", "acospi",
+                 "asin", "asinh", "asinpi",
+                 "atan", "atanh", "atanpi",
+                 "cbrt", "ceil", "cos",
+                 "cosh", "cospi",
+                 "erfc", "erf",
+                 "exp", "exp2", "exp10",
+                 "expm1", "fabs", "floor", "lgamma",
+                 "log", "log2", "log10", "log1p", "logb",
+                 "rint", "round", "rsqrt",
+                 "sin", "sinh", "sinpi",
+                 "sqrt",
+                 "tan", "tanh", "tanpi",
+                 "tgamma", "trunc"] in {
+  def : ConstOCLSPVBuiltin<name, [FGenTypeN, FGenTypeN]>;
+}
+
+foreach name  = ["fmax", "fmin", "fmod",
+                 "atan2", "atan2pi",
+                 "copysign", "fdim", "hypot",
+                 "maxmag", "minmag", "nextafter",
+                 "pow", "powr", "remainder"] in {
+  def : ConstOCLSPVBuiltin<name, [FGenTypeN, FGenTypeN, FGenTypeN]>;
+}
+
+foreach name = ["fma", "mad"] in {
+  def : ConstOCLSPVBuiltin<name, [FGenTypeN, FGenTypeN, FGenTypeN, FGenTypeN]>;
+}
+
+foreach AS = [GlobalAS, LocalAS, PrivateAS, GenericAS, DefaultAS] in {
+  foreach name = ["fract", "modf"] in {
+    def : OCLSPVBuiltin<name, [FGenTypeN, FGenTypeN, PointerType<FGenTypeN, AS>]>;
+  }
+
+  foreach name = ["frexp", "lgamma_r"] in {
+    foreach Type = [GenTypeFloatVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeHalfVecAndScalar] in {
+      def : OCLSPVBuiltin<name, [Type, Type, PointerType<GenTypeIntVecAndScalar, AS>]>;
+    }
+  }
+}
+
+foreach name = ["ilogb"] in {
+  def : ConstOCLSPVBuiltin<name, [GenTypeIntVecAndScalar, GenTypeFloatVecAndScalar]>;
+  def : ConstOCLSPVBuiltin<name, [GenTypeIntVecAndScalar, GenTypeDoubleVecAndScalar]>;
+  def : ConstOCLSPVBuiltin<name, [GenTypeIntVecAndScalar, GenTypeHalfVecAndScalar]>;
+}
+
+foreach name = ["ldexp"] in {
+  foreach Type = [GenTypeFloatVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeHalfVecAndScalar] in {
+    def : ConstOCLSPVBuiltin<name, [Type, Type, GenTypeIntVecAndScalar]>;
+    def : ConstOCLSPVBuiltin<name, [Type, Type, GenTypeUIntVecAndScalar]>;
+  }
+}
+
+foreach name = ["nan"] in {
+  def : ConstOCLSPVBuiltin<name, [GenTypeHalfVecAndScalar, GenTypeShortVecAndScalar]>;
+  def : ConstOCLSPVBuiltin<name, [GenTypeHalfVecAndScalar, GenTypeUShortVecAndScalar]>;
+  def : ConstOCLSPVBuiltin<name, [GenTypeFloatVecAndScalar, GenTypeIntVecAndScalar]>;
+  def : ConstOCLSPVBuiltin<name, [GenTypeFloatVecAndScalar, GenTypeUIntVecAndScalar]>;
+  def : ConstOCLSPVBuiltin<name, [GenTypeDoubleVecAndScalar, GenTypeLongVecAndScalar]>;
+  def : ConstOCLSPVBuiltin<name, [GenTypeDoubleVecAndScalar, GenTypeULongVecAndScalar]>;
+}
+
+foreach name = ["pown"] in {
+  def : ConstOCLSPVBuiltin<name, [GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar, GenTypeIntVecAndScalar]>;
+  def : ConstOCLSPVBuiltin<name, [GenTypeDoubleVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeIntVecAndScalar]>;
+  def : ConstOCLSPVBuiltin<name, [GenTypeHalfVecAndScalar, GenTypeHalfVecAndScalar, GenTypeIntVecAndScalar]>;
+}
+
+foreach AS = [GlobalAS, LocalAS, PrivateAS, GenericAS, DefaultAS] in {
+  foreach name = ["remquo"] in {
+    foreach Type = [GenTypeFloatVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeHalfVecAndScalar] in {
+      def : OCLSPVBuiltin<name, [Type, Type, Type, PointerType<GenTypeIntVecAndScalar, AS>]>;
+    }
+  }
+}
+
+foreach name = ["rootn"] in {
+  def : ConstOCLSPVBuiltin<name, [GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar, GenTypeIntVecAndScalar]>;
+  def : ConstOCLSPVBuiltin<name, [GenTypeDoubleVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeIntVecAndScalar]>;
+  def : ConstOCLSPVBuiltin<name, [GenTypeHalfVecAndScalar, GenTypeHalfVecAndScalar, GenTypeIntVecAndScalar]>;
+}
+
+foreach AS = [GlobalAS, LocalAS, PrivateAS, GenericAS, DefaultAS] in {
+  foreach name = ["sincos"] in {
+    def : OCLSPVBuiltin<name, [FGenTypeN, FGenTypeN, PointerType<FGenTypeN, AS>]>;
+  }
+}
+
+foreach name = ["half_cos",
+                "half_exp", "half_exp2", "half_exp10",
+                "half_log", "half_log2", "half_log10",
+                "half_recip", "half_rsqrt",
+                "half_sin", "half_sqrt", "half_tan",
+                "native_cos", "native_exp", "native_exp2", "native_exp10",
+                "native_log", "native_log2", "native_log10",
+                "native_recip", "native_rsqrt",
+                "native_sin", "native_sqrt", "native_tan"] in {
+  def : ConstOCLSPVBuiltin<name, [GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar]>;
+}
+
+foreach name = ["half_divide", "half_powr", "native_divide", "native_powr"] in {
+  def : ConstOCLSPVBuiltin<name, [GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar]>;
+}
+
+// 2.2. Integer instructions
+
+foreach name = ["clz", "ctz", "popcount"] in {
+  def : ConstOCLSPVBuiltin<name, [AIGenTypeN, AIGenTypeN]>;
+}
+
+def : ConstOCLSPVBuiltin<"rotate", [AIGenTypeN, AIGenTypeN, AIGenTypeN]>;
+
+def : ConstOCLSPVBuiltin<"s_abs", [AUIGenTypeN, ASIGenTypeN]>;
+
+def : ConstOCLSPVBuiltin<"s_abs_diff", [AUIGenTypeN, ASIGenTypeN, ASIGenTypeN]>;
+
+foreach name = ["s_add_sat",
+                "s_hadd", "s_rhadd",
+                "s_max", "s_min",
+                "s_mul_hi", "s_sub_sat"] in {
+  def : ConstOCLSPVBuiltin<name, [ASIGenTypeN, ASIGenTypeN, ASIGenTypeN]>;
+}
+
+foreach name = ["s_clamp", "s_mad_hi", "s_mad_sat"] in {
+ def : ConstOCLSPVBuiltin<name, [ASIGenTypeN, ASIGenTypeN, ASIGenTypeN, ASIGenTypeN]>;
+}
+
+foreach name = ["s_upsample"] in {
+  def : ConstOCLSPVBuiltin<name, [GenTypeShortVecAndScalar, GenTypeCharVecAndScalar, GenTypeUCharVecAndScalar]>;
+  def : ConstOCLSPVBuiltin<name, [GenTypeShortVecAndScalar, GenTypeSCharVecAndScalar, GenTypeUCharVecAndScalar]>;
+  def : ConstOCLSPVBuiltin<name, [GenTypeIntVecAndScalar, GenTypeShortVecAndScalar, GenTypeUShortVecAndScalar]>;
+  def : ConstOCLSPVBuiltin<name, [GenTypeLongVecAndScalar, GenTypeIntVecAndScalar, GenTypeUIntVecAndScalar]>;
+}
+
+def : ConstOCLSPVBuiltin<"s_mad24", [GenTypeIntVecAndScalar, GenTypeIntVecAndScalar, GenTypeIntVecAndScalar, GenTypeIntVecAndScalar]>;
+
+def : ConstOCLSPVBuiltin<"s_mul24", [GenTypeIntVecAndScalar, GenTypeIntVecAndScalar, GenTypeIntVecAndScalar]>;
+
+foreach name = ["u_add_sat", "u_hadd",
+                "u_rhadd",
+                "u_max", "u_min", "u_sub_sat",
+                "u_abs_diff", "u_mul_hi"] in {
+  def : ConstOCLSPVBuiltin<name, [AUIGenTypeN, AUIGenTypeN, AUIGenTypeN]>;
+}
+
+foreach name = ["u_clamp", "u_mad_sat", "u_mad_hi"] in {
+  def : ConstOCLSPVBuiltin<name, [AUIGenTypeN, AUIGenTypeN, AUIGenTypeN, AUIGenTypeN]>;
+}
+
+foreach name = ["u_upsample"] in {
+  def : ConstOCLSPVBuiltin<name, [GenTypeUShortVecAndScalar, GenTypeUCharVecAndScalar, GenTypeUCharVecAndScalar]>;
+  def : ConstOCLSPVBuiltin<name, [GenTypeUIntVecAndScalar, GenTypeUShortVecAndScalar, GenTypeUShortVecAndScalar]>;
+  def : ConstOCLSPVBuiltin<name, [GenTypeULongVecAndScalar, GenTypeUIntVecAndScalar, GenTypeUIntVecAndScalar]>;
+}
+
+def : ConstOCLSPVBuiltin<"u_mad24", [GenTypeUIntVecAndScalar, GenTypeUIntVecAndScalar, GenTypeUIntVecAndScalar, GenTypeUIntVecAndScalar]>;
+
+def : ConstOCLSPVBuiltin<"u_mul24", [GenTypeUIntVecAndScalar, GenTypeUIntVecAndScalar, GenTypeUIntVecAndScalar]>;
+
+def : ConstOCLSPVBuiltin<"u_abs", [AUIGenTypeN, AUIGenTypeN]>;
+
+// 2.3. Common instructions
+
+foreach name = ["degrees", "radians", "sign"] in {
+  def : ConstOCLSPVBuiltin<name, [FGenTypeN, FGenTypeN]>;
+}
+
+foreach name = ["fmax_common", "fmin_common", "step"] in {
+  def : ConstOCLSPVBuiltin<name, [FGenTypeN, FGenTypeN, FGenTypeN]>;
+}
+
+foreach name = ["fclamp", "mix", "smoothstep"] in {
+  def : ConstOCLSPVBuiltin<name, [FGenTypeN, FGenTypeN, FGenTypeN, FGenTypeN]>;
+}
+
+// 2.4. Geometric instructions
+
+foreach name = ["cross"] in {
+  foreach VSize = [3, 4] in {
+    def : ConstOCLSPVBuiltin<name, [VectorType<Float, VSize>, VectorType<Float, VSize>, VectorType<Float, VSize>]>;
+    def : ConstOCLSPVBuiltin<name, [VectorType<Double, VSize>, VectorType<Double, VSize>, VectorType<Double, VSize>]>;
+    def : ConstOCLSPVBuiltin<name, [VectorType<Half, VSize>, VectorType<Half, VSize>, VectorType<Half, VSize>]>;
+  }
+}
+
+foreach name = ["distance"] in {
+  def : ConstOCLSPVBuiltin<name, [Float, GenTypeFloatVec1234, GenTypeFloatVec1234]>;
+  def : ConstOCLSPVBuiltin<name, [Double, GenTypeDoubleVec1234, GenTypeDoubleVec1234]>;
+  def : ConstOCLSPVBuiltin<name, [Half, GenTypeHalfVec1234, GenTypeHalfVec1234]>;
+}
+
+foreach name = ["length"] in {
+  def : ConstOCLSPVBuiltin<name, [Float, GenTypeFloatVec1234]>;
+  def : ConstOCLSPVBuiltin<name, [Double, GenTypeDoubleVec1234]>;
+  def : ConstOCLSPVBuiltin<name, [Half, GenTypeHalfVec1234]>;
+}
+
+foreach name = ["normalize"] in {
+  def : ConstOCLSPVBuiltin<name, [GenTypeFloatVec1234, GenTypeFloatVec1234]>;
+  def : ConstOCLSPVBuiltin<name, [GenTypeDoubleVec1234, GenTypeDoubleVec1234]>;
+  def : ConstOCLSPVBuiltin<name, [GenTypeHalfVec1234, GenTypeHalfVec1234]>;
+}
+
+def : ConstOCLSPVBuiltin<"fast_distance", [Float, GenTypeFloatVec1234, GenTypeFloatVec1234]>;
+
+def : ConstOCLSPVBuiltin<"fast_length", [Float, GenTypeFloatVec1234]>;
+
+def : ConstOCLSPVBuiltin<"fast_normalize", [GenTypeFloatVec1234, GenTypeFloatVec1234]>;
+
+// 2.5. Relational instructions
+
+def : ConstOCLSPVBuiltin<"bitselect", [AGenTypeN, AGenTypeN, AGenTypeN, AGenTypeN]>;
+
+foreach name = ["select"] in {
+  def : ConstOCLSPVBuiltin<name, [SGenTypeN, SGenTypeN, SGenTypeN, SGenTypeN]>;
+  def : ConstOCLSPVBuiltin<name, [SGenTypeN, SGenTypeN, SGenTypeN, UGenTypeN]>;
+  def : ConstOCLSPVBuiltin<name, [UGenTypeN, UGenTypeN, UGenTypeN, UGenTypeN]>;
+  def : ConstOCLSPVBuiltin<name, [UGenTypeN, UGenTypeN, UGenTypeN, SGenTypeN]>;
+  def : ConstOCLSPVBuiltin<name, [GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar, GenTypeUIntVecAndScalar]>;
+  def : ConstOCLSPVBuiltin<name, [GenTypeDoubleVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeULongVecAndScalar]>;
+  def : ConstOCLSPVBuiltin<name, [GenTypeHalfVecAndScalar, GenTypeHalfVecAndScalar, GenTypeHalfVecAndScalar, GenTypeUShortVecAndScalar]>;
+  def : ConstOCLSPVBuiltin<name, [GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar, GenTypeIntVecAndScalar]>;
+  def : ConstOCLSPVBuiltin<name, [GenTypeDoubleVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeLongVecAndScalar]>;
+  def : ConstOCLSPVBuiltin<name, [GenTypeHalfVecAndScalar, GenTypeHalfVecAndScalar, GenTypeHalfVecAndScalar, GenTypeShortVecAndScalar]>;
+}
+
+// 2.6. Vector Data Load and Store instructions
+
+foreach VSize = [2, 3, 4, 8, 16] in {
+  foreach AS = [GlobalAS, LocalAS, PrivateAS, ConstantAS, GenericAS, DefaultAS] in {
+    foreach Ty = TLAll.List in {
+      foreach name = ["vloadn"] in {
+        def : OCLSPVBuiltin<name # "_R" # Ty.Name # VSize, [VectorType<Ty, VSize>, Size, PointerType<ConstType<Ty>, AS>]>;
+      }
+    }
+    foreach name = ["vloada_halfn", "vload_halfn"] in {
+      def : OCLSPVBuiltin<name # "_Rfloat" # VSize, [VectorType<Float, VSize>, Size, PointerType<ConstType<Half>, AS>]>;
+    }
+  }
+  foreach AS = [GlobalAS, LocalAS, PrivateAS, GenericAS, DefaultAS] in {
+    foreach Ty = TLAll.List in {
+      foreach name = ["vstoren"] in {
+        def : OCLSPVBuiltin<name, [Void, VectorType<Ty, VSize>, Size, PointerType<Ty, AS>]>;
+      }
+    }
+    foreach name = ["vstore_halfn", "vstorea_halfn"] in {
+      def : OCLSPVBuiltin<name, [Void, VectorType<Float, VSize>, Size, PointerType<Half, AS>]>;
+      def : OCLSPVBuiltin<name, [Void, VectorType<Double, VSize>, Size, PointerType<Half, AS>]>;
+    }
+    foreach name = ["vstore_halfn_r", "vstorea_halfn_r"] in {
+      def : OCLSPVBuiltin<name, [Void, VectorType<Float, VSize>, Size, PointerType<Half, AS>, UInt]>;
+      def : OCLSPVBuiltin<name, [Void, VectorType<Double, VSize>, Size, PointerType<Half, AS>, UInt]>;
+    }
+  }
+}
+
+foreach AS = [GlobalAS, LocalAS, PrivateAS, ConstantAS, GenericAS, DefaultAS] in {
+  foreach name = ["vload_half"] in {
+    def : OCLSPVBuiltin<name, [Float, Size, PointerType<ConstType<Half>, AS>]>;
+  }
+}
+
+foreach AS = [GlobalAS, LocalAS, PrivateAS, GenericAS, DefaultAS] in {
+  foreach name = ["vstore_half"] in {
+    def : OCLSPVBuiltin<name, [Void, Float, Size, PointerType<Half, AS>]>;
+    def : OCLSPVBuiltin<name, [Void, Double, Size, PointerType<Half, AS>]>;
+  }
+  foreach name = ["vstore_half_r"] in {
+    def : OCLSPVBuiltin<name, [Void, Float, Size, PointerType<Half, AS>, UInt]>;
+    def : OCLSPVBuiltin<name, [Void, Double, Size, PointerType<Half, AS>, UInt]>;
+  }
+}
+
+// 2.7. Miscellaneous Vector instructions
+
+foreach VSize1 = [Vec2, Vec4, Vec8, Vec16] in {
+  foreach VSize2 = [Vec2, Vec4, Vec8, Vec16] in {
+    def : OCLSPVBuiltin<"shuffle", [GenericType<"TLAll" # VSize1.Name, TLAll, VSize1>,
+                              GenericType<"TLAll" # VSize2.Name, TLAll, VSize2>,
+                              GenericType<"TLAllUnsigned" # VSize1.Name, TLAllUnsigned, VSize1>],
+                  Attr.Const>;
+  }
+}
+foreach VSize1 = [Vec2, Vec4, Vec8, Vec16] in {
+  foreach VSize2 = [Vec2, Vec4, Vec8, Vec16] in {
+    def : OCLSPVBuiltin<"shuffle2", [GenericType<"TLAll" # VSize1.Name, TLAll, VSize1>,
+                               GenericType<"TLAll" # VSize2.Name, TLAll, VSize2>,
+                               GenericType<"TLAll" # VSize2.Name, TLAll, VSize2>,
+                               GenericType<"TLAllUnsigned" # VSize1.Name, TLAllUnsigned, VSize1>],
+                  Attr.Const>;
+  }
+}
+
+// 2.8. Misc instructions
+
+let IsVariadic = 1 in {
+  foreach name = ["printf"] in {
+    def : OCLSPVBuiltin<name, [Int, PointerType<ConstType<Char>, ConstantAS>]>;
+  }
+}
+
+foreach name = ["prefetch"] in {
+  def : OCLSPVBuiltin<name, [Void, PointerType<ConstType<AGenTypeN>, GlobalAS>, Size]>;
+}
+
+
+// Core builtins
+
+// 3.32.8. Memory Instructions
+
+foreach name = ["GenericPtrMemSemantics"] in {
+  def : SPVBuiltin<name, [Int, PointerType<ConstType<Int>, GenericAS>], Attr.Const>;
+}
+
+// 3.32.11. Conversion Instructions
+foreach rnd = ["", "_rte", "_rtn", "_rtp", "_rtz"] in {
+  foreach IType = TLUnsignedInts.List in {
+    foreach FType = TLFloat.List in {
+      foreach sat = ["", "_sat"] in {
+        def : SPVBuiltin<"ConvertFToU_R" # IType.Name # sat # rnd, [IType, FType], Attr.Const>;
+      }
+      def : SPVBuiltin<"ConvertUToF_R" # FType.Name # rnd, [FType, IType], Attr.Const>;
+      foreach v = [2, 3, 4, 8, 16] in {
+        foreach sat = ["", "_sat"] in {
+          def : SPVBuiltin<"ConvertFToU_R" # IType.Name # v # sat # rnd,
+                           [VectorType<IType, v>, VectorType<FType, v>],
+                           Attr.Const>;
+        }
+        def : SPVBuiltin<"ConvertUToF_R" # FType.Name # v # rnd,
+                         [VectorType<FType, v>, VectorType<IType, v>],
+                         Attr.Const>;
+      }
+    }
+  }
+
+  foreach IType = TLSignedInts.List in {
+    foreach FType = TLFloat.List in {
+      foreach sat = ["", "_sat"] in {
+        def : SPVBuiltin<"ConvertFToS_R" # IType.Name # sat # rnd, [IType, FType], Attr.Const>;
+      }
+      def : SPVBuiltin<"ConvertSToF_R" # FType.Name # rnd, [FType, IType], Attr.Const>;
+      foreach v = [2, 3, 4, 8, 16] in {
+        foreach sat = ["", "_sat"] in {
+          def : SPVBuiltin<"ConvertFToS_R" # IType.Name # v # sat # rnd,
+                           [VectorType<IType, v>, VectorType<FType, v>],
+                           Attr.Const>;
+        }
+        def : SPVBuiltin<"ConvertSToF_R" # FType.Name # v # rnd,
+                         [VectorType<FType, v>, VectorType<IType, v>],
+                         Attr.Const>;
+      }
+    }
+  }
+
+  foreach InType = TLFloat.List in {
+    foreach OutType = TLFloat.List in {
+      if !ne(OutType.ElementSize, InType.ElementSize) then {
+        def : SPVBuiltin<"FConvert_R" # OutType.Name # rnd, [OutType, InType], Attr.Const>;
+        foreach v = [2, 3, 4, 8, 16] in {
+          def : SPVBuiltin<"FConvert_R" # OutType.Name # v # rnd,
+                           [VectorType<OutType, v>, VectorType<InType, v>],
+                           Attr.Const>;
+        }
+      }
+    }
+  }
+}
+
+foreach sat = ["", "_sat"] in {
+  foreach InType = TLAllInts.List in {
+    foreach OutType = TLUnsignedInts.List in {
+      if !ne(OutType.ElementSize, InType.ElementSize) then {
+        def : SPVBuiltin<"UConvert_R" # OutType.Name # sat, [OutType, InType], Attr.Const>;
+        foreach v = [2, 3, 4, 8, 16] in {
+          def : SPVBuiltin<"UConvert_R" # OutType.Name # v # sat,
+                           [VectorType<OutType, v>, VectorType<InType, v>],
+                           Attr.Const>;
+        }
+      }
+    }
+    foreach OutType = TLSignedInts.List in {
+      if !ne(OutType.ElementSize, InType.ElementSize) then {
+        def : SPVBuiltin<"SConvert_R" # OutType.Name # sat, [OutType, InType], Attr.Const>;
+        foreach v = [2, 3, 4, 8, 16] in {
+          def : SPVBuiltin<"SConvert_R" # OutType.Name # v # sat,
+                           [VectorType<OutType, v>, VectorType<InType, v>],
+                           Attr.Const>;
+        }
+      }
+    }
+  }
+}
+
+foreach InType = TLSignedInts.List in {
+  foreach OutType = TLUnsignedInts.List in {
+    def : SPVBuiltin<"SatConvertSToU_R" # OutType.Name, [OutType, InType], Attr.Const>;
+    foreach v = [2, 3, 4, 8, 16] in {
+      def : SPVBuiltin<"SatConvertSToU_R" # OutType.Name # v,
+                       [VectorType<OutType, v>, VectorType<InType, v>],
+                       Attr.Const>;
+    }
+  }
+}
+
+foreach InType = TLUnsignedInts.List in {
+  foreach OutType = TLSignedInts.List in {
+    def : SPVBuiltin<"SatConvertUToS_R" # OutType.Name, [OutType, InType], Attr.Const>;
+    foreach v = [2, 3, 4, 8, 16] in {
+      def : SPVBuiltin<"SatConvertUToS_R" # OutType.Name # v,
+                       [VectorType<OutType, v>, VectorType<InType, v>],
+                       Attr.Const>;
+    }
+  }
+}
+
+foreach AS = [GlobalAS, LocalAS, PrivateAS] in {
+  def : SPVBuiltin<"GenericCastToPtrExplicit", [PointerType<Char, AS>, PointerType<Char, GenericAS>], Attr.Const>;
+}
+
+foreach Type = TLFloat.List in {
+  foreach v = [2, 3, 4, 8, 16] in {
+    def : SPVBuiltin<"VectorTimesScalar", [VectorType<Type, v>, VectorType<Type, v>, Type], Attr.Const>;
+  }
+}
+
+foreach name = ["Dot"] in {
+  def : SPVBuiltin<name, [Float, GenTypeFloatVecNoScalar, GenTypeFloatVecNoScalar], Attr.Const>;
+  def : SPVBuiltin<name, [Double, GenTypeDoubleVecNoScalar, GenTypeDoubleVecNoScalar], Attr.Const>;
+  def : SPVBuiltin<name, [Half, GenTypeHalfVecNoScalar, GenTypeHalfVecNoScalar], Attr.Const>;
+}
+
+foreach name = ["Any", "All"] in {
+  def : SPVBuiltin<name, [Bool, GenTypeSCharVecNoScalar], Attr.Const>;
+}
+
+foreach name = ["IsNan", "IsInf", "IsFinite", "IsNormal", "SignBitSet"] in {
+  def : SPVBuiltin<name, [Bool, Float], Attr.Const>;
+  def : SPVBuiltin<name, [Bool, Double], Attr.Const>;
+  def : SPVBuiltin<name, [Bool, Half], Attr.Const>;
+  def : SPVBuiltin<name, [GenTypeSCharVecNoScalar, GenTypeFloatVecNoScalar], Attr.Const>;
+  def : SPVBuiltin<name, [GenTypeSCharVecNoScalar, GenTypeDoubleVecNoScalar], Attr.Const>;
+  def : SPVBuiltin<name, [GenTypeSCharVecNoScalar, GenTypeHalfVecNoScalar], Attr.Const>;
+}
+
+foreach name = ["LessOrGreater",
+                "Ordered", "Unordered",
+                "FOrdEqual", "FUnordEqual",
+                "FOrdNotEqual", "FUnordNotEqual",
+                "FOrdLessThan", "FUnordLessThan",
+                "FOrdGreaterThan", "FUnordGreaterThan",
+                "FOrdLessThanEqual", "FUnordLessThanEqual",
+                "FOrdGreaterThanEqual", "FUnordGreaterThanEqual"] in {
+  def : SPVBuiltin<name, [Bool, Float, Float], Attr.Const>;
+  def : SPVBuiltin<name, [Bool, Double, Double], Attr.Const>;
+  def : SPVBuiltin<name, [Bool, Half, Half], Attr.Const>;
+  def : SPVBuiltin<name, [GenTypeSCharVecNoScalar, GenTypeFloatVecNoScalar, GenTypeFloatVecNoScalar], Attr.Const>;
+  def : SPVBuiltin<name, [GenTypeSCharVecNoScalar, GenTypeDoubleVecNoScalar, GenTypeDoubleVecNoScalar], Attr.Const>;
+  def : SPVBuiltin<name, [GenTypeSCharVecNoScalar, GenTypeHalfVecNoScalar, GenTypeHalfVecNoScalar], Attr.Const>;
+}
+
+foreach name = ["BitCount"] in {
+  def : SPVBuiltin<name, [AIGenTypeN, AIGenTypeN], Attr.Const>;
+}
+
+// 3.32.20. Barrier Instructions
+
+foreach name = ["ControlBarrier"] in {
+  // TODO: Allow enum flags instead of UInt ?
+  // TODO: We should enforce that the UInt must be a literal.
+  def : SPVBuiltin<name, [Void, UInt, UInt, UInt], Attr.Convergent>;
+}
+
+foreach name = ["MemoryBarrier"] in {
+  // TODO: Allow enum flags instead of UInt ?
+  // TODO: We should enforce that the UInt must be a literal.
+  def : SPVBuiltin<name, [Void, UInt, UInt]>;
+}
+
+// 3.32.21. Group and Subgroup Instructions
+
+foreach name = ["GroupAsyncCopy"] in {
+  // TODO: Allow enum flags instead of UInt ?
+  // TODO: We should enforce that the UInt must be a literal.
+  def : SPVBuiltin<name, [Event, UInt, PointerType<AGenTypeN, LocalAS>, PointerType<ConstType<AGenTypeN>, GlobalAS>, Size, Size, Event], Attr.Convergent>;
+  def : SPVBuiltin<name, [Event, UInt, PointerType<AGenTypeN, GlobalAS>, PointerType<ConstType<AGenTypeN>, LocalAS>, Size, Size, Event], Attr.Convergent>;
+}
+
+foreach name = ["GroupWaitEvents"] in {
+  def : SPVBuiltin<name, [Void, UInt, Int, PointerType<Event, DefaultAS>], Attr.Convergent>;
+  def : SPVBuiltin<name, [Void, UInt, Int, PointerType<Event, PrivateAS>], Attr.Convergent>;
+  def : SPVBuiltin<name, [Void, UInt, Int, PointerType<Event, GenericAS>], Attr.Convergent>;
+}
+
+foreach name = ["GroupAll", "GroupAny"] in {
+  def : SPVBuiltin<name, [Bool, UInt, Bool], Attr.Convergent>;
+}
+
+foreach name = ["GroupBroadcast"] in {
+  foreach IDType = TLAllInts.List in {
+    def : SPVBuiltin<name, [AGenType1, UInt, AGenType1, IDType], Attr.Convergent>;
+    def : SPVBuiltin<name, [AGenType1, UInt, AGenType1, VectorType<IDType, 2>], Attr.Convergent>;
+    def : SPVBuiltin<name, [AGenType1, UInt, AGenType1, VectorType<IDType, 3>], Attr.Convergent>;
+    def : SPVBuiltin<name, [Bool, UInt, Bool, IDType], Attr.Convergent>;
+    def : SPVBuiltin<name, [Bool, UInt, Bool, VectorType<IDType, 2>], Attr.Convergent>;
+    def : SPVBuiltin<name, [Bool, UInt, Bool, VectorType<IDType, 3>], Attr.Convergent>;
+  }
+}
+
+foreach name = ["GroupIAdd", "GroupNonUniformIMul", "GroupNonUniformBitwiseOr",
+                "GroupNonUniformBitwiseXor", "GroupNonUniformBitwiseAnd"] in {
+  def : SPVBuiltin<name, [AIGenTypeN, UInt, UInt, AIGenTypeN], Attr.Convergent>;
+}
+
+foreach name = ["GroupFAdd", "GroupFMin", "GroupFMax",
+                "GroupNonUniformFMul"] in {
+  def : SPVBuiltin<name, [FGenTypeN, UInt, UInt, FGenTypeN], Attr.Convergent>;
+}
+
+foreach name = ["GroupUMin", "GroupUMax"] in {
+  def : SPVBuiltin<name, [AUIGenTypeN, UInt, UInt, AUIGenTypeN], Attr.Convergent>;
+}
+
+foreach name = ["GroupSMin", "GroupSMax"] in {
+  def : SPVBuiltin<name, [ASIGenTypeN, UInt, UInt, ASIGenTypeN], Attr.Convergent>;
+}

clang/lib/Sema/Sema.cpp

if (getLangOpts().MSVCCompat) {
     if (getLangOpts().CPlusPlus &&
         IdResolver.begin(&Context.Idents.get("type_info")) == IdResolver.end())
       PushOnScopeChains(Context.buildImplicitRecord("type_info", TTK_Class),
                         TUScope);
 
     addImplicitTypedef("size_t", Context.getSizeType());
   }
 
+  if (getLangOpts().SYCLIsDevice) {
+    addImplicitTypedef("__ocl_event_t", Context.OCLEventTy);
+    addImplicitTypedef("__ocl_sampler_t", Context.OCLSamplerTy);
+#ifdef SEMA_STRINGIZE
+#error "Undefine SEMA_STRINGIZE macro."
+#endif
+#define SEMA_STRINGIZE(s) #s
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  addImplicitTypedef(SEMA_STRINGIZE(__ocl_##ImgType##_##Suffix##_t),           \
+                     Context.SingletonId);
+#include "clang/Basic/OpenCLImageTypes.def"
+#undef SEMA_STRINGIZE
+  }
+
+  if (getLangOpts().SYCLIsDevice || getLangOpts().OpenCL) {
+#ifdef SEMA_STRINGIZE
+#error "Undefine SEMA_STRINGIZE macro."
+#endif
+#define SEMA_STRINGIZE(s) #s
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  addImplicitTypedef(SEMA_STRINGIZE(__ocl_sampled_##ImgType##_##Suffix##_t),   \
+                     Context.Sampled##SingletonId);
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
+#undef SEMA_STRINGIZE
+  }
+
   // Initialize predefined OpenCL types and supported extensions and (optional)
   // core features.
   if (getLangOpts().OpenCL) {
     getOpenCLOptions().addSupport(
         Context.getTargetInfo().getSupportedOpenCLOpts(), getLangOpts());
     addImplicitTypedef("sampler_t", Context.OCLSamplerTy);
     addImplicitTypedef("event_t", Context.OCLEventTy);
     if (getLangOpts().OpenCLCPlusPlus || getLangOpts().OpenCLVersion >= 200) {

clang/lib/Sema/SemaDecl.cpp

 /// Returns true if there hasn't been any invalid type diagnosed.
 static bool diagnoseOpenCLTypes(Sema &Se, VarDecl *NewVD) {
   DeclContext *DC = NewVD->getDeclContext();
   QualType R = NewVD->getType();
 
   // OpenCL v2.0 s6.9.b - Image type can only be used as a function argument.
   // OpenCL v2.0 s6.13.16.1 - Pipe type can only be used as a function
   // argument.
-  if (R->isImageType() || R->isPipeType()) {
+  if (!R->isSampledImageType() && (R->isImageType() || R->isPipeType())) {
     Se.Diag(NewVD->getLocation(),
             diag::err_opencl_type_can_only_be_used_as_function_parameter)
         << R;
     NewVD->setInvalidDecl();
     return false;
   }
 
   // OpenCL v1.2 s6.9.r:

clang/lib/Sema/SemaExpr.cpp

static bool isPlaceholderToRemoveAsArg(QualType type) {
   const BuiltinType *placeholder = dyn_cast<BuiltinType>(type);
   if (!placeholder) return false;
 
   switch (placeholder->getKind()) {
   // Ignore all the non-placeholder types.
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
   case BuiltinType::Id:
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  case BuiltinType::Sampled##Id:
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
 #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \
   case BuiltinType::Id:
 #include "clang/Basic/OpenCLExtensionTypes.def"
   // In practice we'll never use this, since all SVE types are sugared
   // via TypedefTypes rather than exposed directly as BuiltinTypes.
 #define SVE_TYPE(Name, Id, SingletonId) \
   case BuiltinType::Id:
 #include "clang/Basic/AArch64SVEACLETypes.def"
ExprResult Sema::CheckPlaceholderExpr(Expr *E) {
 
   case BuiltinType::OMPIterator:
     return ExprError(Diag(E->getBeginLoc(), diag::err_omp_iterator_use));
 
   // Everything else should be impossible.
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
   case BuiltinType::Id:
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  case BuiltinType::Sampled##Id:
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
 #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \
   case BuiltinType::Id:
 #include "clang/Basic/OpenCLExtensionTypes.def"
 #define SVE_TYPE(Name, Id, SingletonId) \
   case BuiltinType::Id:
 #include "clang/Basic/AArch64SVEACLETypes.def"
 #define PPC_VECTOR_TYPE(Name, Id, Size) \
   case BuiltinType::Id:

clang/lib/Sema/SemaLookup.cpp

 #include "llvm/Support/ErrorHandling.h"
 #include <algorithm>
 #include <iterator>
 #include <list>
 #include <set>
 #include <utility>
 #include <vector>
 
+static inline clang::QualType GetFloat16Type(clang::ASTContext &Context);
+
 #include "OpenCLBuiltins.inc"
+#include "SPIRVBuiltins.inc"
 
 using namespace clang;
 using namespace sema;
 
 namespace {
   class UnqualUsingEntry {
     const DeclContext *Nominated;
     const DeclContext *CommonAncestor;

 
 LLVM_DUMP_METHOD void LookupResult::dump() {
   llvm::errs() << "lookup results for " << getLookupName().getAsString()
                << ":\n";
   for (NamedDecl *D : *this)
     D->dump();
 }
 
+static inline QualType GetFloat16Type(clang::ASTContext &Context) {
+  return Context.getLangOpts().OpenCL ? Context.HalfTy : Context.Float16Ty;
+}
+
 /// Diagnose a missing builtin type.
 static QualType diagOpenCLBuiltinTypeError(Sema &S, llvm::StringRef TypeClass,
                                            llvm::StringRef Name) {
   S.Diag(SourceLocation(), diag::err_opencl_type_not_found)
       << TypeClass << Name;
   return S.Context.VoidTy;
 }
 
static QualType getOpenCLTypedefType(Sema &S, llvm::StringRef Name) {
   if (Result.empty())
     return diagOpenCLBuiltinTypeError(S, "typedef", Name);
   TypedefNameDecl *Decl = Result.getAsSingle<TypedefNameDecl>();
   if (!Decl)
     return diagOpenCLBuiltinTypeError(S, "typedef", Name);
   return S.Context.getTypedefType(Decl);
 }
 
-/// Get the QualType instances of the return type and arguments for an OpenCL
+/// Get the QualType instances of the return type and arguments for a ProgModel
 /// builtin function signature.
-/// \param S (in) The Sema instance.
-/// \param OpenCLBuiltin (in) The signature currently handled.
+/// \param Context (in) The Context instance.
+/// \param Builtin (in) The signature currently handled.
 /// \param GenTypeMaxCnt (out) Maximum number of types contained in a generic
 ///        type used as return type or as argument.
 ///        Only meaningful for generic types, otherwise equals 1.
 /// \param RetTypes (out) List of the possible return types.
 /// \param ArgTypes (out) List of the possible argument types.  For each
 ///        argument, ArgTypes contains QualTypes for the Cartesian product
 ///        of (vector sizes) x (types) .
-static void GetQualTypesForOpenCLBuiltin(
-    Sema &S, const OpenCLBuiltinStruct &OpenCLBuiltin, unsigned &GenTypeMaxCnt,
-    SmallVector<QualType, 1> &RetTypes,
+template <typename ProgModel>
+static void GetQualTypesForProgModelBuiltin(
+    Sema &S, const typename ProgModel::BuiltinStruct &Builtin,
+    unsigned &GenTypeMaxCnt, SmallVector<QualType, 1> &RetTypes,
     SmallVector<SmallVector<QualType, 1>, 5> &ArgTypes) {
   // Get the QualType instances of the return types.
-  unsigned Sig = SignatureTable[OpenCLBuiltin.SigTableIndex];
-  OCL2Qual(S, TypeTable[Sig], RetTypes);
+  unsigned Sig = ProgModel::SignatureTable[Builtin.SigTableIndex];
+  ProgModel::Bultin2Qual(S, ProgModel::TypeTable[Sig], RetTypes);
   GenTypeMaxCnt = RetTypes.size();
 
   // Get the QualType instances of the arguments.
   // First type is the return type, skip it.
-  for (unsigned Index = 1; Index < OpenCLBuiltin.NumTypes; Index++) {
+  for (unsigned Index = 1; Index < Builtin.NumTypes; Index++) {
     SmallVector<QualType, 1> Ty;
-    OCL2Qual(S, TypeTable[SignatureTable[OpenCLBuiltin.SigTableIndex + Index]],
-             Ty);
+    ProgModel::Bultin2Qual(
+        S,
+        ProgModel::TypeTable[ProgModel::SignatureTable[Builtin.SigTableIndex +
+                                                       Index]],
     GenTypeMaxCnt = (Ty.size() > GenTypeMaxCnt) ? Ty.size() : GenTypeMaxCnt;
     ArgTypes.push_back(std::move(Ty));
   }
 }
 
-/// Create a list of the candidate function overloads for an OpenCL builtin
+/// Create a list of the candidate function overloads for a ProgModel builtin
 /// function.
 /// \param Context (in) The ASTContext instance.
 /// \param GenTypeMaxCnt (in) Maximum number of types contained in a generic
 ///        type used as return type or as argument.
 ///        Only meaningful for generic types, otherwise equals 1.
 /// \param FunctionList (out) List of FunctionTypes.
 /// \param RetTypes (in) List of the possible return types.
 /// \param ArgTypes (in) List of the possible types for the arguments.
-static void GetOpenCLBuiltinFctOverloads(
+static void GetProgModelBuiltinFctOverloads(
     ASTContext &Context, unsigned GenTypeMaxCnt,
     std::vector<QualType> &FunctionList, SmallVector<QualType, 1> &RetTypes,
-    SmallVector<SmallVector<QualType, 1>, 5> &ArgTypes) {
+    SmallVector<SmallVector<QualType, 1>, 5> &ArgTypes, bool IsVariadic) {
   FunctionProtoType::ExtProtoInfo PI(
       Context.getDefaultCallingConvention(false, false, true));
-  PI.Variadic = false;
+  PI.Variadic = IsVariadic;
 
   // Do not attempt to create any FunctionTypes if there are no return types,
   // which happens when a type belongs to a disabled extension.
   if (RetTypes.size() == 0)
     return;
 
   // Create FunctionTypes for each (gen)type.
   for (unsigned IGenType = 0; IGenType < GenTypeMaxCnt; IGenType++) {
for (unsigned A = 0; A < ArgTypes.size(); A++) {
       ArgList.push_back(ArgTypes[A][Idx]);
     }
 
     FunctionList.push_back(Context.getFunctionType(
         RetTypes[(RetTypes.size() != 1) ? IGenType : 0], ArgList, PI));
   }
 }
 
-/// When trying to resolve a function name, if isOpenCLBuiltin() returns a
-/// non-null <Index, Len> pair, then the name is referencing an OpenCL
+template <typename ProgModel>
+static bool isVersionInMask(const LangOptions &O, unsigned Mask);
+template <>
+bool isVersionInMask<OpenCLBuiltin>(const LangOptions &LO, unsigned Mask) {
+  return isOpenCLVersionContainedInMask(LO, Mask);
+}
+
+// SPIRV Builtins are always permitted, since all builtins are 'SPIRV_ALL'. We
+// have no corresponding language option to check, so we always include them.
+template <>
+bool isVersionInMask<SPIRVBuiltin>(const LangOptions &LO, unsigned Mask) {
+  return true;
+}
+
+/// When trying to resolve a function name, if ProgModel::isBuiltin() returns a
+/// non-null <Index, Len> pair, then the name is referencing a
 /// builtin function.  Add all candidate signatures to the LookUpResult.
 ///
 /// \param S (in) The Sema instance.
 /// \param LR (inout) The LookupResult instance.
 /// \param II (in) The identifier being resolved.
 /// \param FctIndex (in) Starting index in the BuiltinTable.
 /// \param Len (in) The signature list has Len elements.
-static void InsertOCLBuiltinDeclarationsFromTable(Sema &S, LookupResult &LR,
-                                                  IdentifierInfo *II,
-                                                  const unsigned FctIndex,
-                                                  const unsigned Len) {
+template <typename ProgModel>
+static void InsertBuiltinDeclarationsFromTable(
+    Sema &S, LookupResult &LR, IdentifierInfo *II, const unsigned FctIndex,
+    const unsigned Len,
+    std::function<void(const typename ProgModel::BuiltinStruct &,
+                       FunctionDecl &)>
+        ProgModelFinalizer) {
   // The builtin function declaration uses generic types (gentype).
   bool HasGenType = false;
 
   // Maximum number of types contained in a generic type used as return type or
   // as argument.  Only meaningful for generic types, otherwise equals 1.
   unsigned GenTypeMaxCnt;
 
   ASTContext &Context = S.Context;
 
   for (unsigned SignatureIndex = 0; SignatureIndex < Len; SignatureIndex++) {
-    const OpenCLBuiltinStruct &OpenCLBuiltin =
-        BuiltinTable[FctIndex + SignatureIndex];
+    const typename ProgModel::BuiltinStruct &Builtin =
+        ProgModel::BuiltinTable[FctIndex + SignatureIndex];
 
     // Ignore this builtin function if it is not available in the currently
     // selected language version.
-    if (!isOpenCLVersionContainedInMask(Context.getLangOpts(),
-                                        OpenCLBuiltin.Versions))
+    if (!isVersionInMask<ProgModel>(Context.getLangOpts(), Builtin.Versions))
       continue;
 
     // Ignore this builtin function if it carries an extension macro that is
     // not defined. This indicates that the extension is not supported by the
     // target, so the builtin function should not be available.
-    StringRef Extensions = FunctionExtensionTable[OpenCLBuiltin.Extension];
+    StringRef Extensions = ProgModel::FunctionExtensionTable[Builtin.Extension];
     if (!Extensions.empty()) {
       SmallVector<StringRef, 2> ExtVec;
       Extensions.split(ExtVec, " ");
       bool AllExtensionsDefined = true;
       for (StringRef Ext : ExtVec) {
         if (!S.getPreprocessor().isMacroDefined(Ext)) {
           AllExtensionsDefined = false;
           break;
         }
       }
       if (!AllExtensionsDefined)
         continue;
     }
 
     SmallVector<QualType, 1> RetTypes;
     SmallVector<SmallVector<QualType, 1>, 5> ArgTypes;
 
     // Obtain QualType lists for the function signature.
-    GetQualTypesForOpenCLBuiltin(S, OpenCLBuiltin, GenTypeMaxCnt, RetTypes,
-                                 ArgTypes);
+    GetQualTypesForProgModelBuiltin<ProgModel>(S, Builtin, GenTypeMaxCnt,
+                                               RetTypes, ArgTypes);
     if (GenTypeMaxCnt > 1) {
       HasGenType = true;
     }
 
     // Create function overload for each type combination.
     std::vector<QualType> FunctionList;
-    GetOpenCLBuiltinFctOverloads(Context, GenTypeMaxCnt, FunctionList, RetTypes,
-                                 ArgTypes);
+    GetProgModelBuiltinFctOverloads(Context, GenTypeMaxCnt, FunctionList,
+                                    RetTypes, ArgTypes, Builtin.IsVariadic);
 
     SourceLocation Loc = LR.getNameLoc();
     DeclContext *Parent = Context.getTranslationUnitDecl();
-    FunctionDecl *NewOpenCLBuiltin;
+    FunctionDecl *NewBuiltin;
 
     for (const auto &FTy : FunctionList) {
-      NewOpenCLBuiltin = FunctionDecl::Create(
-          Context, Parent, Loc, Loc, II, FTy, /*TInfo=*/nullptr, SC_Extern,
-          S.getCurFPFeatures().isFPConstrained(), false,
-          FTy->isFunctionProtoType());
-      NewOpenCLBuiltin->setImplicit();
+      NewBuiltin = FunctionDecl::Create(Context, Parent, Loc, Loc, II, FTy,
+                                        /*TInfo=*/nullptr, SC_Extern,
+                                        S.getCurFPFeatures().isFPConstrained(),
+                                        false, FTy->isFunctionProtoType());
+      NewBuiltin->setImplicit();
 
       // Create Decl objects for each parameter, adding them to the
       // FunctionDecl.
       const auto *FP = cast<FunctionProtoType>(FTy);
       SmallVector<ParmVarDecl *, 4> ParmList;
       for (unsigned IParm = 0, e = FP->getNumParams(); IParm != e; ++IParm) {
         ParmVarDecl *Parm = ParmVarDecl::Create(
-            Context, NewOpenCLBuiltin, SourceLocation(), SourceLocation(),
-            nullptr, FP->getParamType(IParm), nullptr, SC_None, nullptr);
+            Context, NewBuiltin, SourceLocation(), SourceLocation(), nullptr,
+            FP->getParamType(IParm), nullptr, SC_None, nullptr);
         Parm->setScopeInfo(0, IParm);
         ParmList.push_back(Parm);
       }
-      NewOpenCLBuiltin->setParams(ParmList);
+      NewBuiltin->setParams(ParmList);
 
       // Add function attributes.
-      if (OpenCLBuiltin.IsPure)
-        NewOpenCLBuiltin->addAttr(PureAttr::CreateImplicit(Context));
-      if (OpenCLBuiltin.IsConst)
-        NewOpenCLBuiltin->addAttr(ConstAttr::CreateImplicit(Context));
-      if (OpenCLBuiltin.IsConv)
-        NewOpenCLBuiltin->addAttr(ConvergentAttr::CreateImplicit(Context));
-
+      if (Builtin.IsPure)
+        NewBuiltin->addAttr(PureAttr::CreateImplicit(Context));
+      if (Builtin.IsConst)
+        NewBuiltin->addAttr(ConstAttr::CreateImplicit(Context));
+      if (Builtin.IsConv)
+        NewBuiltin->addAttr(ConvergentAttr::CreateImplicit(Context));
       if (!S.getLangOpts().OpenCLCPlusPlus)
-        NewOpenCLBuiltin->addAttr(OverloadableAttr::CreateImplicit(Context));
+        NewBuiltin->addAttr(OverloadableAttr::CreateImplicit(Context));
 
-      LR.addDecl(NewOpenCLBuiltin);
+      ProgModelFinalizer(Builtin, *NewBuiltin);
+      LR.addDecl(NewBuiltin);
     }
   }
 
   // If we added overloads, need to resolve the lookup result.
   if (Len > 1 || HasGenType)
     LR.resolveKind();
 }
 
if (II) {
         } else if (II == getASTContext().getTypePackElementName()) {
           R.addDecl(getASTContext().getTypePackElementDecl());
           return true;
         }
       }
 
       // Check if this is an OpenCL Builtin, and if so, insert its overloads.
       if (getLangOpts().OpenCL && getLangOpts().DeclareOpenCLBuiltins) {
-        auto Index = isOpenCLBuiltin(II->getName());
+        auto Index = OpenCLBuiltin::isBuiltin(II->getName());
         if (Index.first) {
-          InsertOCLBuiltinDeclarationsFromTable(*this, R, II, Index.first - 1,
-                                                Index.second);
+          InsertBuiltinDeclarationsFromTable<OpenCLBuiltin>(
+              *this, R, II, Index.first - 1, Index.second,
+              [this](const OpenCLBuiltin::BuiltinStruct &OpenCLBuiltin,
+                     FunctionDecl &NewOpenCLBuiltin) {
+                if (!this->getLangOpts().OpenCLCPlusPlus)
+                  NewOpenCLBuiltin.addAttr(
+                      OverloadableAttr::CreateImplicit(Context));
+              });
+          return true;
+        }
+      }
+
+      // Check if this is a SPIR-V Builtin, and if so, insert its overloads.
+      if (getLangOpts().DeclareSPIRVBuiltins) {
+        auto Index = SPIRVBuiltin::isBuiltin(II->getName());
+        if (Index.first) {
+          InsertBuiltinDeclarationsFromTable<SPIRVBuiltin>(
+              *this, R, II, Index.first - 1, Index.second,
+              [this](const SPIRVBuiltin::BuiltinStruct &,
+                     FunctionDecl &NewBuiltin) {
+                if (!this->getLangOpts().CPlusPlus)
+                  NewBuiltin.addAttr(OverloadableAttr::CreateImplicit(Context));
+              });
           return true;
         }
       }
 
       // If this is a builtin on this (or all) targets, create the decl.
       if (unsigned BuiltinID = II->getBuiltinID()) {
         // In C++ and OpenCL (spec v1.2 s6.9.f), we don't have any predefined
         // library functions like 'malloc'. Instead, we'll just error.

clang/lib/Sema/SemaType.cpp

case DeclaratorChunk::Function: {
             diag::err_parameters_retval_cannot_have_fp16_type) << 1;
           D.setInvalidType(true);
         }
       }
 
       if (LangOpts.OpenCL) {
         // OpenCL v2.0 s6.12.5 - A block cannot be the return value of a
         // function.
-        if (T->isBlockPointerType() || T->isImageType() || T->isSamplerT() ||
-            T->isPipeType()) {
+        if (!T->isSampledImageType() &&
+            (T->isBlockPointerType() || T->isImageType() || T->isSamplerT() ||
+             T->isPipeType())) {
           S.Diag(D.getIdentifierLoc(), diag::err_opencl_invalid_return)
               << T << 1 /*hint off*/;
           D.setInvalidType(true);
         }
         // OpenCL doesn't support variadic functions and blocks
         // (s6.9.e and s6.12.5 OpenCL v2.0) except for printf.
         // We also allow here any toolchain reserved identifiers.
         if (FTI.isVariadic &&

clang/lib/Serialization/ASTCommon.cpp

serialization::TypeIdxFromBuiltin(const BuiltinType *BT) {
   case BuiltinType::ObjCSel:
     ID = PREDEF_TYPE_OBJC_SEL;
     break;
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
   case BuiltinType::Id: \
     ID = PREDEF_TYPE_##Id##_ID; \
     break;
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  case BuiltinType::Sampled##Id:                                               \
+    ID = PREDEF_TYPE_SAMPLED_##Id##_ID;                                        \
+    break;
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
 #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \
   case BuiltinType::Id: \
     ID = PREDEF_TYPE_##Id##_ID; \
     break;
 #include "clang/Basic/OpenCLExtensionTypes.def"
   case BuiltinType::OCLSampler:
     ID = PREDEF_TYPE_SAMPLER_ID;
     break;

clang/lib/Serialization/ASTReader.cpp

if (Index < NUM_PREDEF_TYPE_IDS) {
     case PREDEF_TYPE_OBJC_SEL:
       T = Context.ObjCBuiltinSelTy;
       break;
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
     case PREDEF_TYPE_##Id##_ID: \
       T = Context.SingletonId; \
       break;
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  case PREDEF_TYPE_SAMPLED_##Id##_ID:                                          \
+    T = Context.Sampled##SingletonId;                                          \
+    break;
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
 #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \
     case PREDEF_TYPE_##Id##_ID: \
       T = Context.Id##Ty; \
       break;
 #include "clang/Basic/OpenCLExtensionTypes.def"
     case PREDEF_TYPE_SAMPLER_ID:
       T = Context.OCLSamplerTy;
       break;

clang/test/CodeGenOpenCL/sampled_image.cl

+// RUN: %clang_cc1 %s -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 %s -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -o - -cl-std=clc++ | FileCheck %s
+
+__attribute__((overloadable)) void my_read_image(__ocl_sampled_image1d_ro_t img);
+__attribute__((overloadable)) void my_read_image(__ocl_sampled_image2d_ro_t img);
+
+void test_read_image(__ocl_sampled_image1d_ro_t img_ro, __ocl_sampled_image2d_ro_t img_2d) {
+  // CHECK: call void @_Z13my_read_image32__spirv_SampledImage__image1d_ro(%spirv.SampledImage.image1d_ro_t* %{{[0-9]+}})
+  my_read_image(img_ro);
+  // CHECK: call void @_Z13my_read_image32__spirv_SampledImage__image2d_ro(%spirv.SampledImage.image2d_ro_t* %{{[0-9]+}})
+  my_read_image(img_2d);
+}

clang/test/CodeGenSPIRV/spirv-builtin-lookup-win.cpp

+// RUN: %clang_cc1 -triple x86_64-windows-msvc -fdeclare-spirv-builtins -fsyntax-only -emit-llvm %s -o - | FileCheck %s
+
+float acos(float val) {
+  // CHECK: @"?acos@@YAMM@Z"
+  // CHECK: call float @"?__spirv_ocl_acos@@YAMM@Z"
+  return __spirv_ocl_acos(val);
+}
+
+// CHECK: declare dso_local float @"?__spirv_ocl_acos@@YAMM@Z"(float)
+
+double acos(double val) {
+  // CHECK: @"?acos@@YANN@Z"
+  // CHECK: call double @"?__spirv_ocl_acos@@YANN@Z"
+  return __spirv_ocl_acos(val);
+}
+
+// CHECK: declare dso_local double @"?__spirv_ocl_acos@@YANN@Z"(double)

clang/test/CodeGenSPIRV/spirv-builtin-lookup.cpp

+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fdeclare-spirv-builtins -fsyntax-only -emit-llvm %s -o - | FileCheck %s
+
+float acos(float val) {
+  // CHECK: @_Z4acosf
+  // CHECK: call float @_Z16__spirv_ocl_acosf
+  return __spirv_ocl_acos(val);
+}
+
+// CHECK: declare float @_Z16__spirv_ocl_acosf(float)
+
+double acos(double val) {
+  // CHECK: @_Z4acosd
+  // CHECK: call double @_Z16__spirv_ocl_acosd
+  return __spirv_ocl_acos(val);
+}
+
+// CHECK: declare double @_Z16__spirv_ocl_acosd(double)

clang/test/Preprocessor/spirv-macro.cpp

+// RUN: %clang_cc1 %s -E -dM | FileCheck %s
+// RUN: %clang_cc1 %s  -fdeclare-spirv-builtins  -E -dM | FileCheck --check-prefix=CHECK-SPIRV %s
+
+// CHECK-NOT:#define __SPIRV_BUILTIN_DECLARATIONS__
+
+// CHECK-SPIRV:#define __SPIRV_BUILTIN_DECLARATIONS__

clang/test/SemaOpenCL/sampled_image_overload.cl

+// RUN: %clang_cc1 -ast-dump %s | FileCheck %s
+
+void __attribute__((overloadable)) foo(__ocl_sampled_image1d_ro_t);
+void __attribute__((overloadable)) foo(__ocl_sampled_image2d_ro_t);
+
+// CHECK: FunctionDecl {{.*}} <{{.*}}> line:{{.*}} ker 'void (__private __ocl_sampled_image1d_ro_t, __private __ocl_sampled_image2d_ro_t)'
+void kernel ker(__ocl_sampled_image1d_ro_t src1, __ocl_sampled_image2d_ro_t src2) {
+  // CHECK: CallExpr
+  // CHECK-NEXT: ImplicitCastExpr {{.*}} <{{.*}}> 'void (*)(__private __ocl_sampled_image1d_ro_t)'
+  foo(src1);
+  // CHECK: CallExpr
+  // CHECK-NEXT: ImplicitCastExpr {{.*}} <{{.*}}> 'void (*)(__private __ocl_sampled_image2d_ro_t)'
+  foo(src2);
+}

clang/test/SemaSYCL/spirv-builtin-lookup-invalid.cpp

+// RUN: %clang_cc1 -fdeclare-spirv-builtins -fsyntax-only -verify %s
+
+// Verify that invalid call to __spirv_ocl_acos (no viable overloads) get diagnosed
+
+struct InvalidType {};
+void acos(InvalidType Invalid) {
+  __spirv_ocl_acos(Invalid); // expected-error {{no matching function for call to '__spirv_ocl_acos'}}
+  // expected-note@-1 + {{candidate function not viable: no known conversion from}}
+  // too many params
+  __spirv_ocl_acos(42.f, 42.f); // expected-error {{no matching function for call to '__spirv_ocl_acos'}}
+  // expected-note@-1 + {{candidate function not viable: requires 1 argument, but 2 were provided}}
+}

clang/test/SemaSYCL/spirv-builtin-lookup.cpp

+// RUN: %clang_cc1 -fdeclare-spirv-builtins -fsyntax-only -verify %s
+// expected-no-diagnostics
+
+// Verify that __spirv_ocl_acos is recognized as a builtin
+
+float acos(float val) {
+  return __spirv_ocl_acos(val);
+}
+
+double acos(double val) {
+  return __spirv_ocl_acos(val);
+}
+
+typedef int int4 __attribute__((ext_vector_type(4)));
+typedef float float4 __attribute__((ext_vector_type(4)));
+
+int4 ilogb() {
+  float4 f4 = {0.f, 0.f, 0.f, 0.f};
+  int4 i4 = __spirv_ocl_ilogb(f4);
+  return i4;
+}
+
+double sincos(double val, double *res) {
+  return __spirv_ocl_sincos(val, res);
+}
+
+double dot(float4 v1, float4 v2) {
+  return __spirv_Dot(v1, v2);
+}

clang/tools/libclang/CIndex.cpp

bool CursorVisitor::VisitBuiltinTypeLoc(BuiltinTypeLoc TL) {
   switch (TL.getTypePtr()->getKind()) {
 
   case BuiltinType::Void:
   case BuiltinType::NullPtr:
   case BuiltinType::Dependent:
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
   case BuiltinType::Id:
 #include "clang/Basic/OpenCLImageTypes.def"
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  case BuiltinType::Sampled##Id:
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
 #define EXT_OPAQUE_TYPE(ExtTYpe, Id, Ext) case BuiltinType::Id:
 #include "clang/Basic/OpenCLExtensionTypes.def"
   case BuiltinType::OCLSampler:
   case BuiltinType::OCLEvent:
   case BuiltinType::OCLClkEvent:
   case BuiltinType::OCLQueue:
   case BuiltinType::OCLReserveID:
 #define SVE_TYPE(Name, Id, SingletonId) case BuiltinType::Id:

clang/tools/libclang/CXType.cpp

switch (BT->getKind()) {
     BTCASE(NullPtr);
     BTCASE(Overload);
     BTCASE(Dependent);
     BTCASE(ObjCId);
     BTCASE(ObjCClass);
     BTCASE(ObjCSel);
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) BTCASE(Id);
 #include "clang/Basic/OpenCLImageTypes.def"
-#undef IMAGE_TYPE
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix)                   \
+  BTCASE(Sampled##Id);
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
 #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) BTCASE(Id);
 #include "clang/Basic/OpenCLExtensionTypes.def"
     BTCASE(OCLSampler);
     BTCASE(OCLEvent);
     BTCASE(OCLQueue);
     BTCASE(OCLReserveID);
   default:
     return CXType_Unexposed;
switch (K) {
     TKIND(Auto);
     TKIND(Elaborated);
     TKIND(Pipe);
     TKIND(Attributed);
     TKIND(BFloat16);
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) TKIND(Id);
 #include "clang/Basic/OpenCLImageTypes.def"
 #undef IMAGE_TYPE
+#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) TKIND(Sampled##Id);
+#define IMAGE_WRITE_TYPE(Type, Id, Ext)
+#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext)
+#include "clang/Basic/OpenCLImageTypes.def"
+#undef IMAGE_TYPE
 #define EXT_OPAQUE_TYPE(ExtTYpe, Id, Ext) TKIND(Id);
 #include "clang/Basic/OpenCLExtensionTypes.def"
     TKIND(OCLSampler);
     TKIND(OCLEvent);
     TKIND(OCLQueue);
     TKIND(OCLReserveID);
     TKIND(Atomic);
   }

clang/utils/TableGen/CMakeLists.txt

 set(LLVM_LINK_COMPONENTS Support)
 
 add_tablegen(clang-tblgen CLANG
   ASTTableGen.cpp
   ClangASTNodesEmitter.cpp
   ClangASTPropertiesEmitter.cpp
   ClangAttrEmitter.cpp
   ClangCommentCommandInfoEmitter.cpp
   ClangCommentHTMLNamedCharacterReferenceEmitter.cpp
   ClangCommentHTMLTagsEmitter.cpp
   ClangDataCollectorsEmitter.cpp
   ClangDiagnosticsEmitter.cpp
   ClangOpcodesEmitter.cpp
-  ClangOpenCLBuiltinEmitter.cpp
+  ClangProgModelBuiltinEmitter.cpp
   ClangOptionDocEmitter.cpp
   ClangSACheckersEmitter.cpp
   ClangSyntaxEmitter.cpp
   ClangTypeNodesEmitter.cpp
   MveEmitter.cpp
   NeonEmitter.cpp
   RISCVVEmitter.cpp
   SveEmitter.cpp
   TableGen.cpp
   )
 set_target_properties(clang-tblgen PROPERTIES FOLDER "Clang tablegenning")

clang/utils/TableGen/ClangProgModelBuiltinEmitter.cpp

-//===- ClangOpenCLBuiltinEmitter.cpp - Generate Clang OpenCL Builtin handling
+//===- ClangProgModelBuiltinEmitter.cpp - Generate Clang Builtin handling
 //
 //                     The LLVM Compiler Infrastructure
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
-// These backends consume the definitions of OpenCL builtin functions in
-// clang/lib/Sema/OpenCLBuiltins.td and produce builtin handling code for
+// These backends consume the definitions of builtin functions in
+// clang/lib/Sema/*Builtins.td and produce builtin handling code for
 // inclusion in SemaLookup.cpp, or a test file that calls all declared builtins.
 //
 //===----------------------------------------------------------------------===//
 
 #include "TableGenBackends.h"
 #include "llvm/ADT/MapVector.h"
 #include "llvm/ADT/STLExtras.h"
 #include "llvm/ADT/SmallString.h"

 namespace {
 
 // A list of signatures that are shared by one or more builtin functions.
 struct BuiltinTableEntries {
   SmallVector<StringRef, 4> Names;
   std::vector<std::pair<const Record *, unsigned>> Signatures;
 };
 
-// This tablegen backend emits code for checking whether a function is an
-// OpenCL builtin function. If so, all overloads of this function are
-// added to the LookupResult. The generated include file is used by
-// SemaLookup.cpp
+// This tablegen backend emits code for checking whether a function is a
+// builtin function of a programming model. If so, all overloads of this
+// function are added to the LookupResult. The generated include file is
+// used by SemaLookup.cpp
 //
-// For a successful lookup of e.g. the "cos" builtin, isOpenCLBuiltin("cos")
+// For a successful lookup of e.g. the "cos" builtin, isBuiltin("cos")
 // returns a pair <Index, Len>.
 // BuiltinTable[Index] to BuiltinTable[Index + Len] contains the pairs
 // <SigIndex, SigLen> of the overloads of "cos".
 // SignatureTable[SigIndex] to SignatureTable[SigIndex + SigLen] contains
 // one of the signatures of "cos". The SignatureTable entry can be
 // referenced by other functions, e.g. "sin", to exploit the fact that
-// many OpenCL builtins share the same signature.
+// many builtins may share the same signature.
 //
 // The file generated by this TableGen emitter contains the following:
 //
 //  * Structs and enums to represent types and function signatures.
 //
 //  * const char *FunctionExtensionTable[]
 //    List of space-separated OpenCL extensions.  A builtin references an
 //    entry in this table when the builtin requires a particular (set of)
 //    extension(s) to be enabled.
 //
-//  * OpenCLTypeStruct TypeTable[]
+//  * ProgModelTypeStruct TypeTable[]
 //    Type information for return types and arguments.
 //
 //  * unsigned SignatureTable[]
 //    A list of types representing function signatures.  Each entry is an index
 //    into the above TypeTable.  Multiple entries following each other form a
 //    signature, where the first entry is the return type and subsequent
 //    entries are the argument types.
 //
-//  * OpenCLBuiltinStruct BuiltinTable[]
-//    Each entry represents one overload of an OpenCL builtin function and
+//  * BuiltinStruct BuiltinTable[]
+//    Each entry represents one overload of a builtin function and
 //    consists of an index into the SignatureTable and the number of arguments.
 //
-//  * std::pair<unsigned, unsigned> isOpenCLBuiltin(llvm::StringRef Name)
-//    Find out whether a string matches an existing OpenCL builtin function
+//  * std::pair<unsigned, unsigned> isBuiltin(llvm::StringRef Name)
+//    Find out whether a string matches an existing builtin function
 //    name and return an index into BuiltinTable and the number of overloads.
 //
-//  * void OCL2Qual(Sema&, OpenCLTypeStruct, std::vector<QualType>&)
-//    Convert an OpenCLTypeStruct type to a list of QualType instances.
-//    One OpenCLTypeStruct can represent multiple types, primarily when using
+//  * void Bultin2Qual(Sema&, ProgModelTypeStruct, std::vector<QualType>&)
+//    Convert an ProgModelTypeStruct type to a list of QualType instances.
+//    One ProgModelTypeStruct can represent multiple types, primarily when using
 //    GenTypes.
 //
 class BuiltinNameEmitter {
 public:
-  BuiltinNameEmitter(RecordKeeper &Records, raw_ostream &OS)
-      : Records(Records), OS(OS) {}
+  BuiltinNameEmitter(RecordKeeper &Records, raw_ostream &OS,
+                     llvm::StringRef Family)
+      : Records(Records), OS(OS), Family(Family),
+        ClassName((Family + "Builtin").str()) {}
 
   // Entrypoint to generate the functions and structures for checking
-  // whether a function is an OpenCL builtin function.
+  // whether a function is a builtin function.
   void Emit();
 
 private:
   // A list of indices into the builtin function table.
   using BuiltinIndexListTy = SmallVector<unsigned, 11>;
 
-  // Contains OpenCL builtin functions and related information, stored as
+  // Contains builtin functions and related information, stored as
   // Record instances. They are coming from the associated TableGen file.
   RecordKeeper &Records;
 
   // The output file.
   raw_ostream &OS;
 
+  // Family for which the builtin are for.
+  llvm::StringRef Family;
+
+  // Class for which the builtin are for.
+  std::string ClassName;
+
   // Helper function for BuiltinNameEmitter::EmitDeclarations.  Generate enum
   // definitions in the Output string parameter, and save their Record instances
   // in the List parameter.
   // \param Types (in) List containing the Types to extract.
   // \param TypesSeen (inout) List containing the Types already extracted.
   // \param Output (out) String containing the enums to emit in the output file.
   // \param List (out) List containing the extracted Types, except the Types in
   //        TypesSeen.
   void ExtractEnumTypes(std::vector<Record *> &Types,
                         StringMap<bool> &TypesSeen, std::string &Output,
                         std::vector<const Record *> &List);
 
   // Emit the enum or struct used in the generated file.
   // Populate the TypeList at the same time.
   void EmitDeclarations();
 
   // Parse the Records generated by TableGen to populate the SignaturesList,
   // FctOverloadMap and TypeMap.
   void GetOverloads();
 
-  // Compare two lists of signatures and check that e.g. the OpenCL version,
+  // Compare two lists of signatures and check that e.g. the version,
   // function attributes, and extension are equal for each signature.
   // \param Candidate (in) Entry in the SignatureListMap to check.
   // \param SignatureList (in) List of signatures of the considered function.
   // \returns true if the two lists of signatures are identical.
   bool CanReuseSignature(
       BuiltinIndexListTy *Candidate,
       std::vector<std::pair<const Record *, unsigned>> &SignatureList);
 
   // Group functions with the same list of signatures by populating the
   // SignatureListMap.
   // Some builtin functions have the same list of signatures, for example the
   // "sin" and "cos" functions. To save space in the BuiltinTable, the
-  // "isOpenCLBuiltin" function will have the same output for these two
+  // "isBuiltin" function will have the same output for these two
   // function names.
   void GroupBySignature();
 
   // Emit the FunctionExtensionTable that lists all function extensions.
   void EmitExtensionTable();
 
-  // Emit the TypeTable containing all types used by OpenCL builtins.
+  // Emit the TypeTable containing all types used by the builtins.
   void EmitTypeTable();
 
   // Emit the SignatureTable. This table contains all the possible signatures.
   // A signature is stored as a list of indexes of the TypeTable.
   // The first index references the return type (mandatory), and the followings
   // reference its arguments.
   // E.g.:
   // 15, 2, 15 can represent a function with the signature:
   // int func(float, int)
   // The "int" type being at the index 15 in the TypeTable.
   void EmitSignatureTable();
 
   // Emit the BuiltinTable table. This table contains all the overloads of
-  // each function, and is a struct OpenCLBuiltinDecl.
+  // each function, and is a struct BuiltinDecl.
   // E.g.:
   // // 891 convert_float2_rtn
   //   { 58, 2, 3, 100, 0 },
   // This means that the signature of this convert_float2_rtn overload has
   // 1 argument (+1 for the return type), stored at index 58 in
   // the SignatureTable.  This prototype requires extension "3" in the
   // FunctionExtensionTable.  The last two values represent the minimum (1.0)
-  // and maximum (0, meaning no max version) OpenCL version in which this
+  // and maximum (0, meaning no max version) version in which this
   // overload is supported.
   void EmitBuiltinTable();
 
-  // Emit a StringMatcher function to check whether a function name is an
-  // OpenCL builtin function name.
+  // Emit a StringMatcher function to check whether a function name is a
+  // builtin function name.
   void EmitStringMatcher();
 
   // Emit a function returning the clang QualType instance associated with
   // the TableGen Record Type.
   void EmitQualTypeFinder();
 
   // Contains a list of the available signatures, without the name of the
   // function. Each pair consists of a signature and a cumulative index.
private:
   // E.g.:  The function cos: (float cos(float), double cos(double), ...)
   //        <"cos", <<ptrToPrototype0, 5>,
   //                 <ptrToPrototype1, 35>,
   //                 <ptrToPrototype2, 79>>
   // ptrToPrototype1 has the following signature: <double, double>
   MapVector<StringRef, std::vector<std::pair<const Record *, unsigned>>>
       FctOverloadMap;
 
-  // Contains the map of OpenCL types to their index in the TypeTable.
+  // Contains the map of types to their index in the TypeTable.
   MapVector<const Record *, unsigned> TypeMap;
 
-  // List of OpenCL function extensions mapping extension strings to
+  // List of function extensions mapping extension strings to
   // an index into the FunctionExtensionTable.
   StringMap<unsigned> FunctionExtensionIndex;
 
-  // List of OpenCL type names in the same order as in enum OpenCLTypeID.
+  // List of type names in the same order as in enum TypeID.
   // This list does not contain generic types.
   std::vector<const Record *> TypeList;
 
   // Same as TypeList, but for generic types only.
   std::vector<const Record *> GenTypeList;
 
   // Map an ordered vector of signatures to their original Record instances,
   // and to a list of function names that share these signatures.
public:
   // Entrypoint to generate the functions for testing all OpenCL builtin
   // functions.
   void emit() override;
 };
 
 } // namespace
 
 void BuiltinNameEmitter::Emit() {
-  emitSourceFileHeader("OpenCL Builtin handling", OS);
+  std::string Banner = (Family + " Builtin handling").str();
+  emitSourceFileHeader(Banner, OS);
 
   OS << "#include \"llvm/ADT/StringRef.h\"\n";
   OS << "using namespace clang;\n\n";
 
   // Emit enums and structs.
   EmitDeclarations();
 
   // Parse the Records to populate the internal lists.

 void BuiltinNameEmitter::ExtractEnumTypes(std::vector<Record *> &Types,
                                           StringMap<bool> &TypesSeen,
                                           std::string &Output,
                                           std::vector<const Record *> &List) {
   raw_string_ostream SS(Output);
 
   for (const auto *T : Types) {
     if (TypesSeen.find(T->getValueAsString("Name")) == TypesSeen.end()) {
-      SS << "  OCLT_" + T->getValueAsString("Name") << ",\n";
+      SS << "  TID_" + T->getValueAsString("Name") << ",\n";
       // Save the type names in the same order as their enum value. Note that
       // the Record can be a VectorType or something else, only the name is
       // important.
       List.push_back(T);
       TypesSeen.insert(std::make_pair(T->getValueAsString("Name"), true));
     }
   }
   SS.flush();
 }
 
 void BuiltinNameEmitter::EmitDeclarations() {
+  OS << "class " << ClassName << " {\n\n"
+     << "public:\n\n";
+
   // Enum of scalar type names (float, int, ...) and generic type sets.
-  OS << "enum OpenCLTypeID {\n";
+  OS << "enum TypeID {\n";
 
   StringMap<bool> TypesSeen;
   std::string GenTypeEnums;
   std::string TypeEnums;
 
   // Extract generic types and non-generic types separately, to keep
   // gentypes at the end of the enum which simplifies the special handling
   // for gentypes in SemaLookup.
   std::vector<Record *> GenTypes =
       Records.getAllDerivedDefinitions("GenericType");
   ExtractEnumTypes(GenTypes, TypesSeen, GenTypeEnums, GenTypeList);
 
   std::vector<Record *> Types = Records.getAllDerivedDefinitions("Type");
   ExtractEnumTypes(Types, TypesSeen, TypeEnums, TypeList);
 
   OS << TypeEnums;
   OS << GenTypeEnums;
   OS << "};\n";
 
   // Structure definitions.
   OS << R"(
 // Image access qualifier.
-enum OpenCLAccessQual : unsigned char {
-  OCLAQ_None,
-  OCLAQ_ReadOnly,
-  OCLAQ_WriteOnly,
-  OCLAQ_ReadWrite
+enum AccessQual : unsigned char {
+  AQ_None,
+  AQ_ReadOnly,
+  AQ_WriteOnly,
+  AQ_ReadWrite
 };
 
 // Represents a return type or argument type.
-struct OpenCLTypeStruct {
+struct ProgModelTypeStruct {
   // A type (e.g. float, int, ...).
-  const OpenCLTypeID ID;
+  const TypeID ID;
   // Vector size (if applicable; 0 for scalars and generic types).
   const unsigned VectorWidth;
   // 0 if the type is not a pointer.
   const bool IsPointer : 1;
   // 0 if the type is not const.
   const bool IsConst : 1;
   // 0 if the type is not volatile.
   const bool IsVolatile : 1;
   // Access qualifier.
-  const OpenCLAccessQual AccessQualifier;
+  const AccessQual AccessQualifier;
   // Address space of the pointer (if applicable).
   const LangAS AS;
 };
 
-// One overload of an OpenCL builtin function.
-struct OpenCLBuiltinStruct {
-  // Index of the signature in the OpenCLTypeStruct table.
+// One overload of a builtin function.
+struct BuiltinStruct {
+  // Index of the signature in the ProgModelTypeStruct table.
   const unsigned SigTableIndex;
   // Entries between index SigTableIndex and (SigTableIndex + NumTypes - 1) in
   // the SignatureTable represent the complete signature.  The first type at
   // index SigTableIndex is the return type.
   const unsigned NumTypes;
   // Function attribute __attribute__((pure))
   const bool IsPure : 1;
   // Function attribute __attribute__((const))
   const bool IsConst : 1;
   // Function attribute __attribute__((convergent))
   const bool IsConv : 1;
+  // 0 if the function is not variadic.
+  const bool IsVariadic : 1;
   // OpenCL extension(s) required for this overload.
   const unsigned short Extension;
   // OpenCL versions in which this overload is available.
   const unsigned short Versions;
 };
 
+static const char *FunctionExtensionTable[];
+static const ProgModelTypeStruct TypeTable[];
+static const unsigned short SignatureTable[];
+static const BuiltinStruct BuiltinTable[];
+
+static std::pair<unsigned, unsigned> isBuiltin(llvm::StringRef Name);
+static void Bultin2Qual(Sema &Sema, const ProgModelTypeStruct &Ty,
+                        llvm::SmallVectorImpl<QualType> &QT);
+
 )";
+
+  OS << "};  // class " << ClassName << "\n";
 }
 
 // Verify that the combination of GenTypes in a signature is supported.
 // To simplify the logic for creating overloads in SemaLookup, only allow
 // a signature to contain different GenTypes if these GenTypes represent
 // the same number of actual scalar or vector types.
 //
 // Exit with a fatal error if an unsupported construct is encountered.
 static void VerifySignature(const std::vector<Record *> &Signature,
                             const Record *BuiltinRec) {
   unsigned GenTypeVecSizes = 1;
   unsigned GenTypeTypes = 1;
 
   for (const auto *T : Signature) {
     // Check all GenericType arguments in this signature.
     if (T->isSubClassOf("GenericType")) {
       // Check number of vector sizes.
       unsigned NVecSizes =
           T->getValueAsDef("VectorList")->getValueAsListOfInts("List").size();
       if (NVecSizes != GenTypeVecSizes && NVecSizes != 1) {
         if (GenTypeVecSizes > 1) {
           // We already saw a gentype with a different number of vector sizes.
-          PrintFatalError(BuiltinRec->getLoc(),
+          PrintFatalError(
+              BuiltinRec->getLoc(),
               "number of vector sizes should be equal or 1 for all gentypes "
               "in a declaration");
         }
         GenTypeVecSizes = NVecSizes;
       }
 
       // Check number of data types.
       unsigned NTypes =
           T->getValueAsDef("TypeList")->getValueAsListOfDefs("List").size();
       if (NTypes != GenTypeTypes && NTypes != 1) {
         if (GenTypeTypes > 1) {
           // We already saw a gentype with a different number of types.
-          PrintFatalError(BuiltinRec->getLoc(),
+          PrintFatalError(
+              BuiltinRec->getLoc(),
               "number of types should be equal or 1 for all gentypes "
               "in a declaration");
         }
         GenTypeTypes = NTypes;
       }
     }
   }
 }
for (const auto *B : Builtins) {
     } else {
       SignIndex = it->second;
     }
     FctOverloadMap[BName].push_back(std::make_pair(B, SignIndex));
   }
 }
 
 void BuiltinNameEmitter::EmitExtensionTable() {
-  OS << "static const char *FunctionExtensionTable[] = {\n";
+  OS << "const char * " << ClassName << "::FunctionExtensionTable[] = {\n";
   unsigned Index = 0;
   std::vector<Record *> FuncExtensions =
       Records.getAllDerivedDefinitions("FunctionExtension");
 
   for (const auto &FE : FuncExtensions) {
     // Emit OpenCL extension table entry.
     OS << "  // " << Index << ": " << FE->getName() << "\n"
        << "  \"" << FE->getValueAsString("ExtName") << "\",\n";
 
     // Record index of this extension.
     FunctionExtensionIndex[FE->getName()] = Index++;
   }
   OS << "};\n\n";
 }
 
 void BuiltinNameEmitter::EmitTypeTable() {
-  OS << "static const OpenCLTypeStruct TypeTable[] = {\n";
+  OS << "const " << ClassName << "::ProgModelTypeStruct " << ClassName
+     << "::TypeTable[] = {\n";
   for (const auto &T : TypeMap) {
     const char *AccessQual =
         StringSwitch<const char *>(T.first->getValueAsString("AccessQualifier"))
-            .Case("RO", "OCLAQ_ReadOnly")
-            .Case("WO", "OCLAQ_WriteOnly")
-            .Case("RW", "OCLAQ_ReadWrite")
-            .Default("OCLAQ_None");
+            .Case("RO", "AQ_ReadOnly")
+            .Case("WO", "AQ_WriteOnly")
+            .Case("RW", "AQ_ReadWrite")
+            .Default("AQ_None");
 
     OS << "  // " << T.second << "\n"
-       << "  {OCLT_" << T.first->getValueAsString("Name") << ", "
+       << "  {TID_" << T.first->getValueAsString("Name") << ", "
        << T.first->getValueAsInt("VecWidth") << ", "
        << T.first->getValueAsBit("IsPointer") << ", "
        << T.first->getValueAsBit("IsConst") << ", "
-       << T.first->getValueAsBit("IsVolatile") << ", "
-       << AccessQual << ", "
+       << T.first->getValueAsBit("IsVolatile") << ", " << AccessQual << ", "
        << T.first->getValueAsString("AddrSpace") << "},\n";
   }
   OS << "};\n\n";
 }
 
 void BuiltinNameEmitter::EmitSignatureTable() {
   // Store a type (e.g. int, float, int2, ...). The type is stored as an index
-  // of a struct OpenCLType table. Multiple entries following each other form a
-  // signature.
-  OS << "static const unsigned short SignatureTable[] = {\n";
+  // of a struct ProgModelTypeStruct table. Multiple entries following each
+  // other form a signature.
+  OS << "const unsigned short " << ClassName << "::SignatureTable[] = {\n";
   for (const auto &P : SignaturesList) {
     OS << "  // " << P.second << "\n  ";
     for (const Record *R : P.first) {
       unsigned Entry = TypeMap.find(R)->second;
       if (Entry > USHRT_MAX) {
         // Report an error when seeing an entry that is too large for the
         // current index type (unsigned short).  When hitting this, the type
         // of SignatureTable will need to be changed.
static unsigned short EncodeVersions(unsigned int MinVersion,
   }
 
   return Encoded;
 }
 
 void BuiltinNameEmitter::EmitBuiltinTable() {
   unsigned Index = 0;
 
-  OS << "static const OpenCLBuiltinStruct BuiltinTable[] = {\n";
+  OS << "const " << ClassName << "::BuiltinStruct " << ClassName
+     << "::BuiltinTable[] = {\n";
   for (const auto &SLM : SignatureListMap) {
 
     OS << "  // " << (Index + 1) << ": ";
     for (const auto &Name : SLM.second.Names) {
       OS << Name << ", ";
     }
     OS << "\n";
 
     for (const auto &Overload : SLM.second.Signatures) {
       StringRef ExtName = Overload.first->getValueAsDef("Extension")->getName();
       unsigned int MinVersion =
           Overload.first->getValueAsDef("MinVersion")->getValueAsInt("ID");
       unsigned int MaxVersion =
           Overload.first->getValueAsDef("MaxVersion")->getValueAsInt("ID");
 
       OS << "  { " << Overload.second << ", "
          << Overload.first->getValueAsListOfDefs("Signature").size() << ", "
          << (Overload.first->getValueAsBit("IsPure")) << ", "
          << (Overload.first->getValueAsBit("IsConst")) << ", "
          << (Overload.first->getValueAsBit("IsConv")) << ", "
+         << (Overload.first->getValueAsBit("IsVariadic")) << ", "
          << FunctionExtensionIndex[ExtName] << ", "
          << EncodeVersions(MinVersion, MaxVersion) << " },\n";
       Index++;
     }
   }
   OS << "};\n\n";
 }
 
 bool BuiltinNameEmitter::CanReuseSignature(
     BuiltinIndexListTy *Candidate,
     std::vector<std::pair<const Record *, unsigned>> &SignatureList) {
   assert(Candidate->size() == SignatureList.size() &&
          "signature lists should have the same size");
 
-  auto &CandidateSigs =
-      SignatureListMap.find(Candidate)->second.Signatures;
+  auto &CandidateSigs = SignatureListMap.find(Candidate)->second.Signatures;
   for (unsigned Index = 0; Index < Candidate->size(); Index++) {
     const Record *Rec = SignatureList[Index].first;
     const Record *Rec2 = CandidateSigs[Index].first;
     if (Rec->getValueAsBit("IsPure") == Rec2->getValueAsBit("IsPure") &&
         Rec->getValueAsBit("IsConst") == Rec2->getValueAsBit("IsConst") &&
         Rec->getValueAsBit("IsConv") == Rec2->getValueAsBit("IsConv") &&
+        Rec->getValueAsBit("IsVariadic") == Rec2->getValueAsBit("IsVariadic") &&
         Rec->getValueAsDef("MinVersion")->getValueAsInt("ID") ==
             Rec2->getValueAsDef("MinVersion")->getValueAsInt("ID") &&
         Rec->getValueAsDef("MaxVersion")->getValueAsInt("ID") ==
             Rec2->getValueAsDef("MaxVersion")->getValueAsInt("ID") &&
         Rec->getValueAsDef("Extension")->getName() ==
             Rec2->getValueAsDef("Extension")->getName()) {
       return true;
     }
for (const auto &FctName : SLM.second.Names) {
       SS.flush();
       ValidBuiltins.push_back(
           StringMatcher::StringPair(std::string(FctName), RetStmt));
     }
     CumulativeIndex += Ovl.size();
   }
 
   OS << R"(
-// Find out whether a string matches an existing OpenCL builtin function name.
+// Find out whether a string matches an existing builtin function name.
 // Returns: A pair <0, 0> if no name matches.
 //          A pair <Index, Len> indexing the BuiltinTable if the name is
-//          matching an OpenCL builtin function.
-static std::pair<unsigned, unsigned> isOpenCLBuiltin(llvm::StringRef Name) {
-
+//          matching a builtin function.
 )";
 
+  OS << "std::pair<unsigned, unsigned> " << ClassName
+     << "::isBuiltin(llvm::StringRef Name) {\n\n";
+
   StringMatcher("Name", ValidBuiltins, OS).Emit(0, true);
 
   OS << "  return std::make_pair(0, 0);\n";
-  OS << "} // isOpenCLBuiltin\n";
+  OS << "} // isBuiltin\n";
 }
 
 void BuiltinNameEmitter::EmitQualTypeFinder() {
   OS << R"(
 
+// Convert an ProgModelTypeStruct type to a list of QualTypes.
 static QualType getOpenCLEnumType(Sema &S, llvm::StringRef Name);
 static QualType getOpenCLTypedefType(Sema &S, llvm::StringRef Name);
 
-// Convert an OpenCLTypeStruct type to a list of QualTypes.
+// Convert an ProgModelTypeStruct type to a list of QualTypes.
 // Generic types represent multiple types and vector sizes, thus a vector
 // is returned. The conversion is done in two steps:
 // Step 1: A switch statement fills a vector with scalar base types for the
 //         Cartesian product of (vector sizes) x (types) for generic types,
 //         or a single scalar type for non generic types.
 // Step 2: Qualifiers and other type properties such as vector size are
 //         applied.
-static void OCL2Qual(Sema &S, const OpenCLTypeStruct &Ty,
-                     llvm::SmallVectorImpl<QualType> &QT) {
+)";
+
+  OS << "void " << ClassName
+     << "::Bultin2Qual(Sema &S, const ProgModelTypeStruct &Ty, "
+        "llvm::SmallVectorImpl<QualType> &QT) {\n";
+
+  OS << R"(
   ASTContext &Context = S.Context;
   // Number of scalar types in the GenType.
   unsigned GenTypeNumTypes;
   // Pointer to the list of vector sizes for the GenType.
   llvm::ArrayRef<unsigned> GenVectorSizes;
 )";
 
   // Generate list of vector sizes for each generic type.
   for (const auto *VectList : Records.getAllDerivedDefinitions("IntList")) {
-    OS << "  constexpr unsigned List"
-       << VectList->getValueAsString("Name") << "[] = {";
+    OS << "  constexpr unsigned List" << VectList->getValueAsString("Name")
+       << "[] = {";
     for (const auto V : VectList->getValueAsListOfInts("List")) {
       OS << V << ", ";
     }
     OS << "};\n";
   }
 
   // Step 1.
   // Start of switch statement over all types.
for (auto *IT : ImageTypes) {
     }
   }
 
   // Emit the cases for the image types.  For an image type name, there are 3
   // corresponding QualTypes ("RO", "WO", "RW").  The "AccessQualifier" field
   // tells which one is needed.  Emit a switch statement that puts the
   // corresponding QualType into "QT".
   for (const auto &ITE : ImageTypesMap) {
-    OS << "    case OCLT_" << ITE.getKey() << ":\n"
+    OS << "    case TID_" << ITE.getKey() << ":\n"
        << "      switch (Ty.AccessQualifier) {\n"
-       << "        case OCLAQ_None:\n"
+       << "        case AQ_None:\n"
        << "          llvm_unreachable(\"Image without access qualifier\");\n";
     for (const auto &Image : ITE.getValue()) {
       OS << StringSwitch<const char *>(
                 Image->getValueAsString("AccessQualifier"))
-                .Case("RO", "        case OCLAQ_ReadOnly:\n")
-                .Case("WO", "        case OCLAQ_WriteOnly:\n")
-                .Case("RW", "        case OCLAQ_ReadWrite:\n")
+                .Case("RO", "        case AQ_ReadOnly:\n")
+                .Case("WO", "        case AQ_WriteOnly:\n")
+                .Case("RW", "        case AQ_ReadWrite:\n")
          << "          QT.push_back("
          << Image->getValueAsDef("QTExpr")->getValueAsString("TypeExpr")
          << ");\n"
          << "          break;\n";
     }
     OS << "      }\n"
        << "      break;\n";
   }
 
   // Switch cases for generic types.
   for (const auto *GenType : Records.getAllDerivedDefinitions("GenericType")) {
-    OS << "    case OCLT_" << GenType->getValueAsString("Name") << ": {\n";
+    OS << "    case TID_" << GenType->getValueAsString("Name") << ": {\n";
 
     // Build the Cartesian product of (vector sizes) x (types).  Only insert
     // the plain scalar types for now; other type information such as vector
     // size and type qualifiers will be added after the switch statement.
     std::vector<Record *> BaseTypes =
         GenType->getValueAsDef("TypeList")->getValueAsListOfDefs("List");
 
     // Collect all QualTypes for a single vector size into TypeList.
     OS << "      SmallVector<QualType, " << BaseTypes.size() << "> TypeList;\n";
     for (const auto *T : BaseTypes) {
       StringRef Ext =
           T->getValueAsDef("Extension")->getValueAsString("ExtName");
       if (!Ext.empty()) {
         OS << "      if (S.getPreprocessor().isMacroDefined(\"" << Ext
            << "\")) {\n  ";
       }
+      if (T->getValueAsDef("QTExpr")->isSubClassOf("QualTypeFromFunction"))
         OS << "      TypeList.push_back("
-         << T->getValueAsDef("QTExpr")->getValueAsString("TypeExpr") << ");\n";
+           << T->getValueAsDef("QTExpr")->getValueAsString("TypeExpr")
+           << "(Context));\n";
+      else
+        OS << "      TypeList.push_back("
+           << T->getValueAsDef("QTExpr")->getValueAsString("TypeExpr")
+           << ");\n";
       if (!Ext.empty()) {
         OS << "      }\n";
       }
     }
     OS << "      GenTypeNumTypes = TypeList.size();\n";
 
     // Duplicate the TypeList for every vector size.
     std::vector<int64_t> VectorList =
if (TypesSeen.find(T->getValueAsString("Name")) != TypesSeen.end())
       continue;
     TypesSeen.insert(std::make_pair(T->getValueAsString("Name"), true));
 
     // Check the Type does not have an "abstract" QualType
     auto QT = T->getValueAsDef("QTExpr");
     if (QT->getValueAsBit("IsAbstract") == 1)
       continue;
     // Emit the cases for non generic, non image types.
-    OS << "    case OCLT_" << T->getValueAsString("Name") << ":\n";
+    OS << "    case TID_" << T->getValueAsString("Name") << ":\n";
 
     StringRef Ext = T->getValueAsDef("Extension")->getValueAsString("ExtName");
     // If this type depends on an extension, ensure the extension macro is
     // defined.
     if (!Ext.empty()) {
       OS << "      if (S.getPreprocessor().isMacroDefined(\"" << Ext
          << "\")) {\n  ";
     }
+    if (QT->isSubClassOf("QualTypeFromFunction"))
+      OS << "      QT.push_back(" << QT->getValueAsString("TypeExpr")
+         << "(Context));\n";
+    else
       OS << "      QT.push_back(" << QT->getValueAsString("TypeExpr") << ");\n";
     if (!Ext.empty()) {
       OS << "      }\n";
     }
     OS << "      break;\n";
   }
 
   // End of switch statement.
   OS << "  } // end of switch (Ty.ID)\n\n";
)";
   if (Ty.IsPointer != 0) {
     for (unsigned Index = 0; Index < QT.size(); Index++) {
       QT[Index] = Context.getAddrSpaceQualType(QT[Index], Ty.AS);
       QT[Index] = Context.getPointerType(QT[Index]);
     }
   }
 )";
 
-  // End of the "OCL2Qual" function.
-  OS << "\n} // OCL2Qual\n";
+  // End of the "Bultin2Qual" function.
+  OS << "\n} // Bultin2Qual\n";
 }
 
 std::string OpenCLBuiltinFileEmitterBase::getTypeString(const Record *Type,
                                                         TypeFlags Flags,
                                                         int VectorSize) const {
   std::string S;
   if (Type->getValueAsBit("IsConst") || Flags.IsConst) {
     S += "const ";
for (const auto *B : Builtins) {
     }
 
     OS << OptionalVersionEndif;
     OS << OptionalExtensionEndif;
   }
 }
 
 void clang::EmitClangOpenCLBuiltins(RecordKeeper &Records, raw_ostream &OS) {
-  BuiltinNameEmitter NameChecker(Records, OS);
+  BuiltinNameEmitter NameChecker(Records, OS, "OpenCL");
+  NameChecker.Emit();
+}
+
+void clang::EmitClangSPIRVBuiltins(RecordKeeper &Records, raw_ostream &OS) {
+  BuiltinNameEmitter NameChecker(Records, OS, "SPIRV");
   NameChecker.Emit();
 }
 
 void clang::EmitClangOpenCLBuiltinTests(RecordKeeper &Records,
                                         raw_ostream &OS) {
   OpenCLBuiltinTestEmitter TestFileGenerator(Records, OS);
   TestFileGenerator.emit();
 }

clang/utils/TableGen/TableGen.cpp

enum ActionType {
   GenClangSyntaxNodeClasses,
   GenClangCommentHTMLTags,
   GenClangCommentHTMLTagsProperties,
   GenClangCommentHTMLNamedCharacterReferences,
   GenClangCommentCommandInfo,
   GenClangCommentCommandList,
   GenClangOpenCLBuiltins,
   GenClangOpenCLBuiltinTests,
+  GenClangSPIRVBuiltins,
   GenArmNeon,
   GenArmFP16,
   GenArmBF16,
   GenArmNeonSema,
   GenArmNeonTest,
   GenArmMveHeader,
   GenArmMveBuiltinDef,
   GenArmMveBuiltinSema,
cl::values(
                    "are used in documentation comments"),
         clEnumValN(GenClangCommentCommandList, "gen-clang-comment-command-list",
                    "Generate list of commands that are used in "
                    "documentation comments"),
         clEnumValN(GenClangOpenCLBuiltins, "gen-clang-opencl-builtins",
                    "Generate OpenCL builtin declaration handlers"),
         clEnumValN(GenClangOpenCLBuiltinTests, "gen-clang-opencl-builtin-tests",
                    "Generate OpenCL builtin declaration tests"),
+        clEnumValN(GenClangSPIRVBuiltins, "gen-clang-spirv-builtins",
+                   "Generate SPIR-V builtin declaration handlers"),
         clEnumValN(GenArmNeon, "gen-arm-neon", "Generate arm_neon.h for clang"),
         clEnumValN(GenArmFP16, "gen-arm-fp16", "Generate arm_fp16.h for clang"),
         clEnumValN(GenArmBF16, "gen-arm-bf16", "Generate arm_bf16.h for clang"),
         clEnumValN(GenArmNeonSema, "gen-arm-neon-sema",
                    "Generate ARM NEON sema support for clang"),
         clEnumValN(GenArmNeonTest, "gen-arm-neon-test",
                    "Generate ARM NEON tests for clang"),
         clEnumValN(GenArmSveHeader, "gen-arm-sve-header",
case GenClangCommentCommandList:
     EmitClangCommentCommandList(Records, OS);
     break;
   case GenClangOpenCLBuiltins:
     EmitClangOpenCLBuiltins(Records, OS);
     break;
   case GenClangOpenCLBuiltinTests:
     EmitClangOpenCLBuiltinTests(Records, OS);
     break;
+  case GenClangSPIRVBuiltins:
+    EmitClangSPIRVBuiltins(Records, OS);
+    break;
   case GenClangSyntaxNodeList:
     EmitClangSyntaxNodeList(Records, OS);
     break;
   case GenClangSyntaxNodeClasses:
     EmitClangSyntaxNodeClasses(Records, OS);
     break;
   case GenArmNeon:
     EmitNeon(Records, OS);

clang/utils/TableGen/TableGenBackends.h

 void EmitClangAttrDocs(llvm::RecordKeeper &Records, llvm::raw_ostream &OS);
 void EmitClangDiagDocs(llvm::RecordKeeper &Records, llvm::raw_ostream &OS);
 void EmitClangOptDocs(llvm::RecordKeeper &Records, llvm::raw_ostream &OS);
 
 void EmitClangOpenCLBuiltins(llvm::RecordKeeper &Records,
                              llvm::raw_ostream &OS);
 void EmitClangOpenCLBuiltinTests(llvm::RecordKeeper &Records,
                                  llvm::raw_ostream &OS);
+void EmitClangSPIRVBuiltins(llvm::RecordKeeper &Records, llvm::raw_ostream &OS);
 
 void EmitClangDataCollectors(llvm::RecordKeeper &Records,
                              llvm::raw_ostream &OS);
 
 void EmitTestPragmaAttributeSupportedAttributes(llvm::RecordKeeper &Records,
                                                 llvm::raw_ostream &OS);
 
 } // end namespace clang
 
 #endif