clang/lib/Sema/OpenCLBuiltins.td

   1 //==--- OpenCLBuiltins.td - OpenCL builtin declarations -------------------===//
   2 //
   3 //                     The LLVM Compiler Infrastructure
   4 //
   5 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
   6 // See https://llvm.org/LICENSE.txt for license information.
   7 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
   8 //
   9 //===----------------------------------------------------------------------===//
  10 //
  11 // This file contains TableGen definitions for OpenCL builtin function
  12 // declarations.  In case of an unresolved function name in OpenCL, Clang will
  13 // check for a function described in this file when -fdeclare-opencl-builtins
  14 // is specified.
  15 //
  16 //===----------------------------------------------------------------------===//
  17
  18 //===----------------------------------------------------------------------===//
  19 //              Definitions of miscellaneous basic entities.
  20 //===----------------------------------------------------------------------===//
  21 // Versions of OpenCL
  22 class Version<int _Version> {
  23   int ID = _Version;
  24 }
  25 def CLAll : Version<  0>;
  26 def CL10  : Version<100>;
  27 def CL11  : Version<110>;
  28 def CL12  : Version<120>;
  29 def CL20  : Version<200>;
  30
  31 // Address spaces
  32 // Pointer types need to be assigned an address space.
  33 class AddressSpace<string _AS> {
  34   string Name = _AS;
  35 }
  36 def DefaultAS    : AddressSpace<"clang::LangAS::Default">;
  37 def PrivateAS    : AddressSpace<"clang::LangAS::opencl_private">;
  38 def GlobalAS     : AddressSpace<"clang::LangAS::opencl_global">;
  39 def ConstantAS   : AddressSpace<"clang::LangAS::opencl_constant">;
  40 def LocalAS      : AddressSpace<"clang::LangAS::opencl_local">;
  41 def GenericAS    : AddressSpace<"clang::LangAS::opencl_generic">;
  42
  43 // OpenCL language extension.
  44 class AbstractExtension<string _Ext> {
  45   // One or more OpenCL extensions, space separated.  Each extension must be
  46   // a valid extension name for the opencl extension pragma.
  47   string ExtName = _Ext;
  48 }
  49
  50 // Extension associated to a builtin function.
  51 class FunctionExtension<string _Ext> : AbstractExtension<_Ext>;
  52
  53 // Extension associated to a type.  This enables implicit conditionalization of
  54 // builtin function overloads containing a type that depends on an extension.
  55 // During overload resolution, when a builtin function overload contains a type
  56 // with a TypeExtension, those overloads are skipped when the extension is
  57 // disabled.
  58 class TypeExtension<string _Ext> : AbstractExtension<_Ext>;
  59
  60 // Concatenate zero or more space-separated extensions in NewExts to Base and
  61 // return the resulting FunctionExtension in ret.
  62 class concatExtension<FunctionExtension Base, string NewExts> {
  63   FunctionExtension ret = FunctionExtension<
  64     !cond(
  65       // Return Base extension if NewExts is empty,
  66       !empty(NewExts) : Base.ExtName,
  67
  68       // otherwise, return NewExts if Base extension is empty,
  69       !empty(Base.ExtName) : NewExts,
  70
  71       // otherwise, concatenate NewExts to Base.
  72       true : Base.ExtName # " " # NewExts
  73     )
  74   >;
  75 }
  76
  77 // TypeExtension definitions.
  78 def NoTypeExt   : TypeExtension<"">;
  79 def Fp16TypeExt : TypeExtension<"cl_khr_fp16">;
  80 def Fp64TypeExt : TypeExtension<"cl_khr_fp64">;
  81 def Atomic64TypeExt : TypeExtension<"cl_khr_int64_base_atomics cl_khr_int64_extended_atomics">;
  82 def AtomicFp64TypeExt : TypeExtension<"cl_khr_int64_base_atomics cl_khr_int64_extended_atomics cl_khr_fp64">;
  83
  84 // FunctionExtension definitions.
  85 def FuncExtNone                          : FunctionExtension<"">;
  86 def FuncExtKhrSubgroups                  : FunctionExtension<"__opencl_subgroup_builtins">;
  87 def FuncExtKhrSubgroupExtendedTypes      : FunctionExtension<"cl_khr_subgroup_extended_types">;
  88 def FuncExtKhrSubgroupNonUniformVote     : FunctionExtension<"cl_khr_subgroup_non_uniform_vote">;
  89 def FuncExtKhrSubgroupBallot             : FunctionExtension<"cl_khr_subgroup_ballot">;
  90 def FuncExtKhrSubgroupNonUniformArithmetic: FunctionExtension<"cl_khr_subgroup_non_uniform_arithmetic">;
  91 def FuncExtKhrSubgroupShuffle            : FunctionExtension<"cl_khr_subgroup_shuffle">;
  92 def FuncExtKhrSubgroupShuffleRelative    : FunctionExtension<"cl_khr_subgroup_shuffle_relative">;
  93 def FuncExtKhrSubgroupClusteredReduce    : FunctionExtension<"cl_khr_subgroup_clustered_reduce">;
  94 def FuncExtKhrExtendedBitOps             : FunctionExtension<"cl_khr_extended_bit_ops">;
  95 def FuncExtKhrGlobalInt32BaseAtomics     : FunctionExtension<"cl_khr_global_int32_base_atomics">;
  96 def FuncExtKhrGlobalInt32ExtendedAtomics : FunctionExtension<"cl_khr_global_int32_extended_atomics">;
  97 def FuncExtKhrLocalInt32BaseAtomics      : FunctionExtension<"cl_khr_local_int32_base_atomics">;
  98 def FuncExtKhrLocalInt32ExtendedAtomics  : FunctionExtension<"cl_khr_local_int32_extended_atomics">;
  99 def FuncExtKhrInt64BaseAtomics           : FunctionExtension<"cl_khr_int64_base_atomics">;
 100 def FuncExtKhrInt64ExtendedAtomics       : FunctionExtension<"cl_khr_int64_extended_atomics">;
 101 def FuncExtKhrMipmapImage                : FunctionExtension<"cl_khr_mipmap_image">;
 102 def FuncExtKhrMipmapImageWrites          : FunctionExtension<"cl_khr_mipmap_image_writes">;
 103 def FuncExtKhrGlMsaaSharing              : FunctionExtension<"cl_khr_gl_msaa_sharing">;
 104
 105 def FuncExtOpenCLCDeviceEnqueue          : FunctionExtension<"__opencl_c_device_enqueue">;
 106 def FuncExtOpenCLCGenericAddressSpace    : FunctionExtension<"__opencl_c_generic_address_space">;
 107 def FuncExtOpenCLCNamedAddressSpaceBuiltins : FunctionExtension<"__opencl_c_named_address_space_builtins">;
 108 def FuncExtOpenCLCPipes                  : FunctionExtension<"__opencl_c_pipes">;
 109 def FuncExtOpenCLCWGCollectiveFunctions  : FunctionExtension<"__opencl_c_work_group_collective_functions">;
 110 def FuncExtOpenCLCReadWriteImages        : FunctionExtension<"__opencl_c_read_write_images">;
 111 def FuncExtFloatAtomicsFp16GlobalASLoadStore  : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_global_atomic_load_store">;
 112 def FuncExtFloatAtomicsFp16LocalASLoadStore   : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_local_atomic_load_store">;
 113 def FuncExtFloatAtomicsFp16GenericASLoadStore : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_global_atomic_load_store __opencl_c_ext_fp16_local_atomic_load_store">;
 114 def FuncExtFloatAtomicsFp16GlobalASAdd        : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_global_atomic_add">;
 115 def FuncExtFloatAtomicsFp32GlobalASAdd        : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp32_global_atomic_add">;
 116 def FuncExtFloatAtomicsFp64GlobalASAdd        : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp64_global_atomic_add">;
 117 def FuncExtFloatAtomicsFp16LocalASAdd         : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_local_atomic_add">;
 118 def FuncExtFloatAtomicsFp32LocalASAdd         : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp32_local_atomic_add">;
 119 def FuncExtFloatAtomicsFp64LocalASAdd         : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp64_local_atomic_add">;
 120 def FuncExtFloatAtomicsFp16GenericASAdd       : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_local_atomic_add __opencl_c_ext_fp16_global_atomic_add">;
 121 def FuncExtFloatAtomicsFp32GenericASAdd       : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp32_local_atomic_add __opencl_c_ext_fp32_global_atomic_add">;
 122 def FuncExtFloatAtomicsFp64GenericASAdd       : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp64_local_atomic_add __opencl_c_ext_fp64_global_atomic_add">;
 123 def FuncExtFloatAtomicsFp16GlobalASMinMax     : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_global_atomic_min_max">;
 124 def FuncExtFloatAtomicsFp32GlobalASMinMax     : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp32_global_atomic_min_max">;
 125 def FuncExtFloatAtomicsFp64GlobalASMinMax     : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp64_global_atomic_min_max">;
 126 def FuncExtFloatAtomicsFp16LocalASMinMax      : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_local_atomic_min_max">;
 127 def FuncExtFloatAtomicsFp32LocalASMinMax      : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp32_local_atomic_min_max">;
 128 def FuncExtFloatAtomicsFp64LocalASMinMax      : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp64_local_atomic_min_max">;
 129 def FuncExtFloatAtomicsFp16GenericASMinMax    : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_local_atomic_min_max __opencl_c_ext_fp16_global_atomic_min_max">;
 130 def FuncExtFloatAtomicsFp32GenericASMinMax    : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp32_local_atomic_min_max __opencl_c_ext_fp32_global_atomic_min_max">;
 131 def FuncExtFloatAtomicsFp64GenericASMinMax    : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp64_local_atomic_min_max __opencl_c_ext_fp64_global_atomic_min_max">;
 132
 133 // Not a real extension, but a workaround to add C++ for OpenCL specific builtins.
 134 def FuncExtOpenCLCxx                     : FunctionExtension<"__cplusplus">;
 135
 136 // Arm extensions.
 137 def ArmIntegerDotProductInt8                   : FunctionExtension<"cl_arm_integer_dot_product_int8">;
 138 def ArmIntegerDotProductAccumulateInt8         : FunctionExtension<"cl_arm_integer_dot_product_accumulate_int8">;
 139 def ArmIntegerDotProductAccumulateInt16        : FunctionExtension<"cl_arm_integer_dot_product_accumulate_int16">;
 140 def ArmIntegerDotProductAccumulateSaturateInt8 : FunctionExtension<"cl_arm_integer_dot_product_accumulate_saturate_int8">;
 141
 142 // Qualified Type.  These map to ASTContext::QualType.
 143 class QualType<string _TypeExpr, bit _IsAbstract=0> {
 144   // Expression to obtain the QualType inside OCL2Qual.
 145   // E.g. TypeExpr="Context.IntTy" for the int type.
 146   string TypeExpr = _TypeExpr;
 147   // Some QualTypes in this file represent an abstract type for which there is
 148   // no corresponding AST QualType, e.g. a GenType or an `image2d_t` type
 149   // without access qualifiers.
 150   bit IsAbstract = _IsAbstract;
 151 }
 152
 153 // List of integers.
 154 class IntList<string _Name, list<int> _List> {
 155   string Name = _Name;
 156   list<int> List = _List;
 157 }
 158
 159 //===----------------------------------------------------------------------===//
 160 //                      OpenCL C classes for types
 161 //===----------------------------------------------------------------------===//
 162 // OpenCL C basic data types (int, float, image2d_t, ...).
 163 // Its child classes can represent concrete types (e.g. VectorType) or
 164 // abstract types (e.g. GenType).
 165 class Type<string _Name, QualType _QTExpr> {
 166   // Name of the Type.
 167   string Name = _Name;
 168   // QualType associated with this type.
 169   QualType QTExpr = _QTExpr;
 170   // Size of the vector (if applicable).
 171   int VecWidth = 1;
 172   // Is a pointer.
 173   bit IsPointer = 0;
 174   // "const" qualifier.
 175   bit IsConst = 0;
 176   // "volatile" qualifier.
 177   bit IsVolatile = 0;
 178   // Access qualifier. Must be one of ("RO", "WO", "RW").
 179   string AccessQualifier = "";
 180   // Address space.
 181   string AddrSpace = DefaultAS.Name;
 182   // Extension that needs to be enabled to expose a builtin that uses this type.
 183   TypeExtension Extension = NoTypeExt;
 184 }
 185
 186 // OpenCL vector types (e.g. int2, int3, int16, float8, ...).
 187 class VectorType<Type _Ty, int _VecWidth> : Type<_Ty.Name, _Ty.QTExpr> {
 188   let VecWidth = _VecWidth;
 189   let AccessQualifier = "";
 190   // Inherited fields
 191   let IsPointer = _Ty.IsPointer;
 192   let IsConst = _Ty.IsConst;
 193   let IsVolatile = _Ty.IsVolatile;
 194   let AddrSpace = _Ty.AddrSpace;
 195   let Extension = _Ty.Extension;
 196 }
 197
 198 // OpenCL pointer types (e.g. int*, float*, ...).
 199 class PointerType<Type _Ty, AddressSpace _AS = DefaultAS> :
 200     Type<_Ty.Name, _Ty.QTExpr> {
 201   let AddrSpace = _AS.Name;
 202   // Inherited fields
 203   let VecWidth = _Ty.VecWidth;
 204   let IsPointer = 1;
 205   let IsConst = _Ty.IsConst;
 206   let IsVolatile = _Ty.IsVolatile;
 207   let AccessQualifier = _Ty.AccessQualifier;
 208   let Extension = _Ty.Extension;
 209 }
 210
 211 // OpenCL const types (e.g. const int).
 212 class ConstType<Type _Ty> : Type<_Ty.Name, _Ty.QTExpr> {
 213   let IsConst = 1;
 214   // Inherited fields
 215   let VecWidth = _Ty.VecWidth;
 216   let IsPointer = _Ty.IsPointer;
 217   let IsVolatile = _Ty.IsVolatile;
 218   let AccessQualifier = _Ty.AccessQualifier;
 219   let AddrSpace = _Ty.AddrSpace;
 220   let Extension = _Ty.Extension;
 221 }
 222
 223 // OpenCL volatile types (e.g. volatile int).
 224 class VolatileType<Type _Ty> : Type<_Ty.Name, _Ty.QTExpr> {
 225   let IsVolatile = 1;
 226   // Inherited fields
 227   let VecWidth = _Ty.VecWidth;
 228   let IsPointer = _Ty.IsPointer;
 229   let IsConst = _Ty.IsConst;
 230   let AccessQualifier = _Ty.AccessQualifier;
 231   let AddrSpace = _Ty.AddrSpace;
 232   let Extension = _Ty.Extension;
 233 }
 234
 235 // OpenCL image types (e.g. image2d).
 236 class ImageType<Type _Ty, string _AccessQualifier> :
 237     Type<_Ty.Name, QualType<_Ty.QTExpr.TypeExpr # _AccessQualifier # "Ty", 0>> {
 238   let VecWidth = 0;
 239   let AccessQualifier = _AccessQualifier;
 240   // Inherited fields
 241   let IsPointer = _Ty.IsPointer;
 242   let IsConst = _Ty.IsConst;
 243   let IsVolatile = _Ty.IsVolatile;
 244   let AddrSpace = _Ty.AddrSpace;
 245   // Add TypeExtensions for writable "image3d_t" and "read_write" image types.
 246   let Extension = !cond(
 247       !and(!eq(_Ty.Name, "image3d_t"), !eq(_AccessQualifier, "WO")) : TypeExtension<"cl_khr_3d_image_writes">,
 248       !and(!eq(_Ty.Name, "image3d_t"), !eq(_AccessQualifier, "RW")) : TypeExtension<"cl_khr_3d_image_writes __opencl_c_read_write_images">,
 249       !or(!eq(_Ty.Name, "image2d_depth_t"), !eq(_Ty.Name, "image2d_array_depth_t")) : TypeExtension<"cl_khr_depth_images">,
 250       !eq(_AccessQualifier, "RW") : TypeExtension<"__opencl_c_read_write_images">,
 251       true : _Ty.Extension);
 252 }
 253
 254 // OpenCL enum type (e.g. memory_scope).
 255 class EnumType<string _Name> :
 256     Type<_Name, QualType<"getOpenCLEnumType(S, \"" # _Name # "\")", 0>> {
 257 }
 258
 259 // OpenCL typedef type (e.g. cl_mem_fence_flags).
 260 class TypedefType<string _Name> :
 261     Type<_Name, QualType<"getOpenCLTypedefType(S, \"" # _Name # "\")", 0>> {
 262 }
 263
 264 // List of Types.
 265 class TypeList<list<Type> _Type> {
 266   list<Type> List = _Type;
 267 }
 268
 269 // A GenericType is an abstract type that defines a set of types as a
 270 // combination of Types and vector sizes.
 271 //
 272 // For example, if TypeList = <int, float> and VectorList = <1, 2, 4>, then it
 273 // represents <int, int2, int4, float, float2, float4>.
 274 //
 275 // Some rules apply when using multiple GenericType arguments in a declaration:
 276 //   1. The number of vector sizes must be equal or 1 for all gentypes in a
 277 //      declaration.
 278 //   2. The number of Types must be equal or 1 for all gentypes in a
 279 //      declaration.
 280 //   3. Generic types are combined by iterating over all generic types at once.
 281 //      For example, for the following GenericTypes
 282 //        GenT1 = GenericType<half, [1, 2]> and
 283 //        GenT2 = GenericType<float, int, [1, 2]>
 284 //      A declaration f(GenT1, GenT2) results in the combinations
 285 //        f(half, float), f(half2, float2), f(half, int), f(half2, int2) .
 286 //   4. "sgentype" from the OpenCL specification is supported by specifying
 287 //      a single vector size.
 288 //      For example, for the following GenericTypes
 289 //        GenT = GenericType<half, int, [1, 2]> and
 290 //        SGenT = GenericType<half, int, [1]>
 291 //      A declaration f(GenT, SGenT) results in the combinations
 292 //        f(half, half), f(half2, half), f(int, int), f(int2, int) .
 293 class GenericType<string _Ty, TypeList _TypeList, IntList _VectorList> :
 294     Type<_Ty, QualType<"null", 1>> {
 295   // Possible element types of the generic type.
 296   TypeList TypeList = _TypeList;
 297   // Possible vector sizes of the types in the TypeList.
 298   IntList VectorList = _VectorList;
 299   // The VecWidth field is ignored for GenericTypes. Use VectorList instead.
 300   let VecWidth = 0;
 301 }
 302
 303 // Builtin function attributes.
 304 def Attr {
 305   list<bit> None = [0, 0, 0];
 306   list<bit> Pure = [1, 0, 0];
 307   list<bit> Const = [0, 1, 0];
 308   list<bit> Convergent = [0, 0, 1];
 309 }
 310
 311 //===----------------------------------------------------------------------===//
 312 //                      OpenCL C class for builtin functions
 313 //===----------------------------------------------------------------------===//
 314 class Builtin<string _Name, list<Type> _Signature, list<bit> _Attributes = Attr.None> {
 315   // Name of the builtin function
 316   string Name = _Name;
 317   // List of types used by the function. The first one is the return type and
 318   // the following are the arguments. The list must have at least one element
 319   // (the return type).
 320   list<Type> Signature = _Signature;
 321   // Function attribute __attribute__((pure))
 322   bit IsPure = _Attributes[0];
 323   // Function attribute __attribute__((const))
 324   bit IsConst = _Attributes[1];
 325   // Function attribute __attribute__((convergent))
 326   bit IsConv = _Attributes[2];
 327   // OpenCL extensions to which the function belongs.
 328   FunctionExtension Extension = FuncExtNone;
 329   // Version of OpenCL from which the function is available (e.g.: CL10).
 330   // MinVersion is inclusive.
 331   Version MinVersion = CL10;
 332   // Version of OpenCL from which the function is not supported anymore.
 333   // MaxVersion is exclusive.
 334   // CLAll makes the function available for all versions.
 335   Version MaxVersion = CLAll;
 336 }
 337
 338 //===----------------------------------------------------------------------===//
 339 //                 Definitions of OpenCL C types
 340 //===----------------------------------------------------------------------===//
 341
 342 // OpenCL v1.0/1.2/2.0 s6.1.1: Built-in Scalar Data Types.
 343 def Bool      : Type<"bool",      QualType<"Context.BoolTy">>;
 344 def Char      : Type<"char",      QualType<"Context.CharTy">>;
 345 def UChar     : Type<"uchar",     QualType<"Context.UnsignedCharTy">>;
 346 def Short     : Type<"short",     QualType<"Context.ShortTy">>;
 347 def UShort    : Type<"ushort",    QualType<"Context.UnsignedShortTy">>;
 348 def Int       : Type<"int",       QualType<"Context.IntTy">>;
 349 def UInt      : Type<"uint",      QualType<"Context.UnsignedIntTy">>;
 350 def Long      : Type<"long",      QualType<"Context.LongTy">>;
 351 def ULong     : Type<"ulong",     QualType<"Context.UnsignedLongTy">>;
 352 def Float     : Type<"float",     QualType<"Context.FloatTy">>;
 353 let Extension = Fp64TypeExt in {
 354   def Double    : Type<"double",    QualType<"Context.DoubleTy">>;
 355 }
 356
 357 // The half type for builtins that require the cl_khr_fp16 extension.
 358 let Extension = Fp16TypeExt in {
 359   def Half      : Type<"half",      QualType<"Context.HalfTy">>;
 360 }
 361
 362 // Without the cl_khr_fp16 extension, the half type can only be used to declare
 363 // a pointer.  Define const and non-const pointer types in all address spaces.
 364 // Use the "__half" alias to allow the TableGen emitter to distinguish the
 365 // (extensionless) pointee type of these pointer-to-half types from the "half"
 366 // type defined above that already carries the cl_khr_fp16 extension.
 367 foreach AS = [PrivateAS, GlobalAS, ConstantAS, LocalAS, GenericAS] in {
 368   def "HalfPtr" # AS      : PointerType<Type<"__half", QualType<"Context.HalfTy">>, AS>;
 369   def "HalfPtrConst" # AS : PointerType<ConstType<Type<"__half", QualType<"Context.HalfTy">>>, AS>;
 370 }
 371
 372 def Size      : Type<"size_t",    QualType<"Context.getSizeType()">>;
 373 def PtrDiff   : Type<"ptrdiff_t", QualType<"Context.getPointerDiffType()">>;
 374 def IntPtr    : Type<"intptr_t",  QualType<"Context.getIntPtrType()">>;
 375 def UIntPtr   : Type<"uintptr_t", QualType<"Context.getUIntPtrType()">>;
 376 def Void      : Type<"void",      QualType<"Context.VoidTy">>;
 377
 378 // OpenCL v1.0/1.2/2.0 s6.1.2: Built-in Vector Data Types.
 379 // Built-in vector data types are created by TableGen's OpenCLBuiltinEmitter.
 380
 381 // OpenCL v1.0/1.2/2.0 s6.1.3: Other Built-in Data Types.
 382 // The image definitions are "abstract".  They should not be used without
 383 // specifying an access qualifier (RO/WO/RW).
 384 def Image1d               : Type<"image1d_t", QualType<"Context.OCLImage1d", 1>>;
 385 def Image2d               : Type<"image2d_t", QualType<"Context.OCLImage2d", 1>>;
 386 def Image3d               : Type<"image3d_t", QualType<"Context.OCLImage3d", 1>>;
 387 def Image1dArray          : Type<"image1d_array_t", QualType<"Context.OCLImage1dArray", 1>>;
 388 def Image1dBuffer         : Type<"image1d_buffer_t", QualType<"Context.OCLImage1dBuffer", 1>>;
 389 def Image2dArray          : Type<"image2d_array_t", QualType<"Context.OCLImage2dArray", 1>>;
 390 def Image2dDepth          : Type<"image2d_depth_t", QualType<"Context.OCLImage2dDepth", 1>>;
 391 def Image2dArrayDepth     : Type<"image2d_array_depth_t", QualType<"Context.OCLImage2dArrayDepth", 1>>;
 392 def Image2dMsaa           : Type<"image2d_msaa_t", QualType<"Context.OCLImage2dMSAA", 1>>;
 393 def Image2dArrayMsaa      : Type<"image2d_array_msaa_t", QualType<"Context.OCLImage2dArrayMSAA", 1>>;
 394 def Image2dMsaaDepth      : Type<"image2d_msaa_depth_t", QualType<"Context.OCLImage2dMSAADepth", 1>>;
 395 def Image2dArrayMsaaDepth : Type<"image2d_array_msaa_depth_t", QualType<"Context.OCLImage2dArrayMSAADepth", 1>>;
 396
 397 def Sampler               : Type<"sampler_t", QualType<"Context.OCLSamplerTy">>;
 398 def ClkEvent              : Type<"clk_event_t", QualType<"Context.OCLClkEventTy">>;
 399 def Event                 : Type<"event_t", QualType<"Context.OCLEventTy">>;
 400 def Queue                 : Type<"queue_t", QualType<"Context.OCLQueueTy">>;
 401 def ReserveId             : Type<"reserve_id_t", QualType<"Context.OCLReserveIDTy">>;
 402 def MemFenceFlags         : TypedefType<"cl_mem_fence_flags">;
 403 def ClkProfilingInfo      : TypedefType<"clk_profiling_info">;
 404 def NDRange               : TypedefType<"ndrange_t">;
 405
 406 // OpenCL v2.0 s6.13.11: Atomic integer and floating-point types.
 407 def AtomicInt             : Type<"atomic_int", QualType<"Context.getAtomicType(Context.IntTy)">>;
 408 def AtomicUInt            : Type<"atomic_uint", QualType<"Context.getAtomicType(Context.UnsignedIntTy)">>;
 409 let Extension = Atomic64TypeExt in {
 410   def AtomicLong            : Type<"atomic_long", QualType<"Context.getAtomicType(Context.LongTy)">>;
 411   def AtomicULong           : Type<"atomic_ulong", QualType<"Context.getAtomicType(Context.UnsignedLongTy)">>;
 412 }
 413 def AtomicFloat           : Type<"atomic_float", QualType<"Context.getAtomicType(Context.FloatTy)">>;
 414 let Extension = AtomicFp64TypeExt in {
 415   def AtomicDouble          : Type<"atomic_double", QualType<"Context.getAtomicType(Context.DoubleTy)">>;
 416 }
 417 def AtomicHalf            : Type<"atomic_half", QualType<"Context.getAtomicType(Context.HalfTy)">>;
 418 def AtomicIntPtr          : Type<"atomic_intptr_t", QualType<"Context.getAtomicType(Context.getIntPtrType())">>;
 419 def AtomicUIntPtr         : Type<"atomic_uintptr_t", QualType<"Context.getAtomicType(Context.getUIntPtrType())">>;
 420 def AtomicSize            : Type<"atomic_size_t", QualType<"Context.getAtomicType(Context.getSizeType())">>;
 421 def AtomicPtrDiff         : Type<"atomic_ptrdiff_t", QualType<"Context.getAtomicType(Context.getPointerDiffType())">>;
 422
 423 def AtomicFlag            : TypedefType<"atomic_flag">;
 424 def MemoryOrder           : EnumType<"memory_order">;
 425 def MemoryScope           : EnumType<"memory_scope">;
 426
 427 //===----------------------------------------------------------------------===//
 428 //                 Definitions of OpenCL gentype variants
 429 //===----------------------------------------------------------------------===//
 430 // The OpenCL specification often uses "gentype" in builtin function
 431 // declarations to indicate that a builtin function is available with various
 432 // argument and return types.  The types represented by "gentype" vary between
 433 // different parts of the specification.  The following definitions capture
 434 // the different type lists for gentypes in different parts of the
 435 // specification.
 436
 437 // Vector width lists.
 438 def VecAndScalar: IntList<"VecAndScalar", [1, 2, 3, 4, 8, 16]>;
 439 def VecNoScalar : IntList<"VecNoScalar", [2, 3, 4, 8, 16]>;
 440 def Vec1        : IntList<"Vec1", [1]>;
 441 def Vec1234     : IntList<"Vec1234", [1, 2, 3, 4]>;
 442
 443 // Type lists.
 444 def TLAll           : TypeList<[Char,  UChar, Short,  UShort, Int,  UInt, Long,  ULong, Float, Double, Half]>;
 445 def TLFloat         : TypeList<[Float, Double, Half]>;
 446 def TLSignedInts    : TypeList<[Char, Short, Int, Long]>;
 447 def TLUnsignedInts  : TypeList<[UChar, UShort, UInt, ULong]>;
 448
 449 def TLIntLongFloats : TypeList<[Int, UInt, Long, ULong, Float, Double, Half]>;
 450
 451 // All unsigned integer types twice, to facilitate unsigned return types for e.g.
 452 // uchar abs(char) and
 453 // uchar abs(uchar).
 454 def TLAllUIntsTwice : TypeList<[UChar, UChar, UShort, UShort, UInt, UInt, ULong, ULong]>;
 455
 456 def TLAllInts       : TypeList<[Char, UChar, Short, UShort, Int, UInt, Long, ULong]>;
 457
 458 // GenType definitions for multiple base types (e.g. all floating point types,
 459 // or all integer types).
 460 // All types
 461 def AGenType1              : GenericType<"AGenType1", TLAll, Vec1>;
 462 def AGenTypeN              : GenericType<"AGenTypeN", TLAll, VecAndScalar>;
 463 def AGenTypeNNoScalar      : GenericType<"AGenTypeNNoScalar", TLAll, VecNoScalar>;
 464 // All integer
 465 def AIGenType1             : GenericType<"AIGenType1", TLAllInts, Vec1>;
 466 def AIGenTypeN             : GenericType<"AIGenTypeN", TLAllInts, VecAndScalar>;
 467 def AIGenTypeNNoScalar     : GenericType<"AIGenTypeNNoScalar", TLAllInts, VecNoScalar>;
 468 // All integer to unsigned
 469 def AI2UGenTypeN           : GenericType<"AI2UGenTypeN", TLAllUIntsTwice, VecAndScalar>;
 470 // Signed integer
 471 def SGenTypeN              : GenericType<"SGenTypeN", TLSignedInts, VecAndScalar>;
 472 // Unsigned integer
 473 def UGenTypeN              : GenericType<"UGenTypeN", TLUnsignedInts, VecAndScalar>;
 474 // Float
 475 def FGenTypeN              : GenericType<"FGenTypeN", TLFloat, VecAndScalar>;
 476 // (u)int, (u)long, and all floats
 477 def IntLongFloatGenType1   : GenericType<"IntLongFloatGenType1", TLIntLongFloats, Vec1>;
 478 // (u)char and (u)short
 479 def CharShortGenType1      : GenericType<"CharShortGenType1",
 480                                  TypeList<[Char, UChar, Short, UShort]>, Vec1>;
 481
 482 // GenType definitions for every single base type (e.g. fp32 only).
 483 // Names are like: GenTypeFloatVecAndScalar.
 484 foreach Type = [Char, UChar, Short, UShort,
 485                 Int, UInt, Long, ULong,
 486                 Float, Double, Half] in {
 487   foreach VecSizes = [VecAndScalar, VecNoScalar] in {
 488     def "GenType" # Type # VecSizes :
 489               GenericType<"GenType" # Type # VecSizes,
 490                           TypeList<[Type]>, VecSizes>;
 491   }
 492 }
 493
 494 // GenType definitions for vec1234.
 495 foreach Type = [Float, Double, Half] in {
 496   def "GenType" # Type # Vec1234 :
 497               GenericType<"GenType" # Type # Vec1234,
 498                           TypeList<[Type]>, Vec1234>;
 499 }
 500
 501
 502 //===----------------------------------------------------------------------===//
 503 //                 Definitions of OpenCL builtin functions
 504 //===----------------------------------------------------------------------===//
 505 //--------------------------------------------------------------------
 506 // OpenCL v1.1/1.2/2.0 s6.2.3 - Explicit conversions.
 507 // OpenCL v2.0 Extensions s5.1.1 and s6.1.1 - Conversions.
 508
 509 // Generate the convert_* builtins functions.
 510 foreach RType = [Float, Double, Half, Char, UChar, Short,
 511                  UShort, Int, UInt, Long, ULong] in {
 512   foreach IType = [Float, Double, Half, Char, UChar, Short,
 513                    UShort, Int, UInt, Long, ULong] in {
 514     // Conversions to integer type have a sat and non-sat variant.
 515     foreach sat = !cond(!eq(RType.Name, "float") : [""],
 516                         !eq(RType.Name, "double") : [""],
 517                         !eq(RType.Name, "half") : [""],
 518                         1 : ["", "_sat"]) in {
 519       foreach rnd = ["", "_rte", "_rtn", "_rtp", "_rtz"] in {
 520         def : Builtin<"convert_" # RType.Name # sat # rnd, [RType, IType],
 521                       Attr.Const>;
 522         foreach v = [2, 3, 4, 8, 16] in {
 523           def : Builtin<"convert_" # RType.Name # v # sat # rnd,
 524                         [VectorType<RType, v>, VectorType<IType, v>],
 525                         Attr.Const>;
 526         }
 527       }
 528     }
 529   }
 530 }
 531
 532 //--------------------------------------------------------------------
 533 // OpenCL v1.1 s6.11.1, v1.2 s6.12.1, v2.0 s6.13.1 - Work-item Functions
 534 // --- Table 7 ---
 535 def : Builtin<"get_work_dim", [UInt], Attr.Const>;
 536 foreach name = ["get_global_size", "get_global_id", "get_local_size",
 537                 "get_local_id", "get_num_groups", "get_group_id",
 538                 "get_global_offset"] in {
 539   def : Builtin<name, [Size, UInt], Attr.Const>;
 540 }
 541
 542 let MinVersion = CL20 in {
 543   def : Builtin<"get_enqueued_local_size", [Size, UInt]>;
 544   foreach name = ["get_global_linear_id", "get_local_linear_id"] in {
 545     def : Builtin<name, [Size]>;
 546   }
 547 }
 548
 549
 550 //--------------------------------------------------------------------
 551 // OpenCL v1.1 s6.11.2, v1.2 s6.12.2, v2.0 s6.13.2 - Math functions
 552 // OpenCL Extension v2.0 s5.1.2 and s6.1.2 - Math Functions
 553 // --- Table 8 ---
 554 // --- 1 argument ---
 555 foreach name = ["acos", "acosh", "acospi",
 556                 "asin", "asinh", "asinpi",
 557                 "atan", "atanh", "atanpi",
 558                 "cbrt", "ceil",
 559                 "cos", "cosh", "cospi",
 560                 "erfc", "erf",
 561                 "exp", "exp2", "exp10", "expm1",
 562                 "fabs", "floor",
 563                 "log", "log2", "log10", "log1p", "logb",
 564                 "rint", "round", "rsqrt",
 565                 "sin", "sinh", "sinpi",
 566                 "tan", "tanh", "tanpi",
 567                 "tgamma", "trunc",
 568                 "lgamma"] in {
 569     def : Builtin<name, [FGenTypeN, FGenTypeN], Attr.Const>;
 570 }
 571
 572 // sqrt is handled in opencl-c-base.h to handle
 573 // -cl-fp32-correctly-rounded-divide-sqrt.
 574
 575 foreach name = ["nan"] in {
 576   def : Builtin<name, [GenTypeFloatVecAndScalar, GenTypeUIntVecAndScalar], Attr.Const>;
 577   def : Builtin<name, [GenTypeDoubleVecAndScalar, GenTypeULongVecAndScalar], Attr.Const>;
 578   def : Builtin<name, [GenTypeHalfVecAndScalar, GenTypeUShortVecAndScalar], Attr.Const>;
 579 }
 580
 581 // --- 2 arguments ---
 582 foreach name = ["atan2", "atan2pi", "copysign", "fdim", "fmod", "hypot",
 583                 "maxmag", "minmag", "nextafter", "pow", "powr",
 584                 "remainder"] in {
 585   def : Builtin<name, [FGenTypeN, FGenTypeN, FGenTypeN], Attr.Const>;
 586 }
 587 foreach name = ["fmax", "fmin"] in {
 588   def : Builtin<name, [FGenTypeN, FGenTypeN, FGenTypeN], Attr.Const>;
 589   def : Builtin<name, [GenTypeFloatVecNoScalar, GenTypeFloatVecNoScalar, Float], Attr.Const>;
 590   def : Builtin<name, [GenTypeDoubleVecNoScalar, GenTypeDoubleVecNoScalar, Double], Attr.Const>;
 591   def : Builtin<name, [GenTypeHalfVecNoScalar, GenTypeHalfVecNoScalar, Half], Attr.Const>;
 592 }
 593 foreach name = ["ilogb"] in {
 594   def : Builtin<name, [GenTypeIntVecAndScalar, GenTypeFloatVecAndScalar], Attr.Const>;
 595   def : Builtin<name, [GenTypeIntVecAndScalar, GenTypeDoubleVecAndScalar], Attr.Const>;
 596   def : Builtin<name, [GenTypeIntVecAndScalar, GenTypeHalfVecAndScalar], Attr.Const>;
 597 }
 598 foreach name = ["ldexp"] in {
 599   def : Builtin<name, [GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar, GenTypeIntVecAndScalar], Attr.Const>;
 600   def : Builtin<name, [GenTypeFloatVecNoScalar, GenTypeFloatVecNoScalar, Int], Attr.Const>;
 601   def : Builtin<name, [GenTypeDoubleVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeIntVecAndScalar], Attr.Const>;
 602   def : Builtin<name, [GenTypeDoubleVecNoScalar, GenTypeDoubleVecNoScalar, Int], Attr.Const>;
 603   def : Builtin<name, [GenTypeHalfVecAndScalar, GenTypeHalfVecAndScalar, GenTypeIntVecAndScalar], Attr.Const>;
 604   def : Builtin<name, [GenTypeHalfVecNoScalar, GenTypeHalfVecNoScalar, Int], Attr.Const>;
 605 }
 606 foreach name = ["pown", "rootn"] in {
 607   def : Builtin<name, [GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar, GenTypeIntVecAndScalar], Attr.Const>;
 608   def : Builtin<name, [GenTypeDoubleVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeIntVecAndScalar], Attr.Const>;
 609   def : Builtin<name, [GenTypeHalfVecAndScalar, GenTypeHalfVecAndScalar, GenTypeIntVecAndScalar], Attr.Const>;
 610 }
 611
 612 // --- 3 arguments ---
 613 foreach name = ["fma", "mad"] in {
 614   def : Builtin<name, [FGenTypeN, FGenTypeN, FGenTypeN, FGenTypeN], Attr.Const>;
 615 }
 616
 617 // The following math builtins take pointer arguments.  Which overloads are
 618 // available depends on whether the generic address space feature is enabled.
 619 multiclass MathWithPointer<list<AddressSpace> addrspaces> {
 620   foreach AS = addrspaces in {
 621     foreach name = ["fract", "modf", "sincos"] in {
 622       def : Builtin<name, [FGenTypeN, FGenTypeN, PointerType<FGenTypeN, AS>]>;
 623     }
 624     foreach name = ["frexp", "lgamma_r"] in {
 625       foreach Type = [GenTypeFloatVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeHalfVecAndScalar] in {
 626         def : Builtin<name, [Type, Type, PointerType<GenTypeIntVecAndScalar, AS>]>;
 627       }
 628     }
 629     foreach name = ["remquo"] in {
 630       foreach Type = [GenTypeFloatVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeHalfVecAndScalar] in {
 631         def : Builtin<name, [Type, Type, Type, PointerType<GenTypeIntVecAndScalar, AS>]>;
 632       }
 633     }
 634   }
 635 }
 636
 637 let Extension = FuncExtOpenCLCNamedAddressSpaceBuiltins in {
 638   defm : MathWithPointer<[GlobalAS, LocalAS, PrivateAS]>;
 639 }
 640 let Extension = FuncExtOpenCLCGenericAddressSpace in {
 641   defm : MathWithPointer<[GenericAS]>;
 642 }
 643
 644 // --- Table 9 ---
 645 foreach name = ["half_cos",
 646                 "half_exp", "half_exp2", "half_exp10",
 647                 "half_log", "half_log2", "half_log10",
 648                 "half_recip", "half_rsqrt",
 649                 "half_sin", "half_sqrt", "half_tan",
 650                 "native_cos",
 651                 "native_exp", "native_exp2", "native_exp10",
 652                 "native_log", "native_log2", "native_log10",
 653                 "native_recip", "native_rsqrt",
 654                 "native_sin", "native_sqrt", "native_tan"] in {
 655   def : Builtin<name, [GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar], Attr.Const>;
 656 }
 657 foreach name = ["half_divide", "half_powr",
 658                 "native_divide", "native_powr"] in {
 659   def : Builtin<name, [GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar], Attr.Const>;
 660 }
 661
 662 //--------------------------------------------------------------------
 663 // OpenCL v1.1 s6.11.3, v1.2 s6.12.3, v2.0 s6.13.3 - Integer Functions
 664 // --- Table 10 ---
 665 // --- 1 argument ---
 666 foreach name = ["abs"] in {
 667   def : Builtin<name, [AI2UGenTypeN, AIGenTypeN], Attr.Const>;
 668 }
 669 def : Builtin<"clz", [AIGenTypeN, AIGenTypeN], Attr.Const>;
 670 let MinVersion = CL12 in {
 671   def : Builtin<"popcount", [AIGenTypeN, AIGenTypeN], Attr.Const>;
 672 }
 673 let MinVersion = CL20 in {
 674   foreach name = ["ctz"] in {
 675     def : Builtin<name, [AIGenTypeN, AIGenTypeN], Attr.Const>;
 676   }
 677 }
 678
 679 // --- 2 arguments ---
 680 foreach name = ["abs_diff"] in {
 681   def : Builtin<name, [AI2UGenTypeN, AIGenTypeN, AIGenTypeN], Attr.Const>;
 682 }
 683 foreach name = ["add_sat", "hadd", "rhadd", "mul_hi", "rotate", "sub_sat"] in {
 684   def : Builtin<name, [AIGenTypeN, AIGenTypeN, AIGenTypeN], Attr.Const>;
 685 }
 686 foreach name = ["max", "min"] in {
 687   def : Builtin<name, [AIGenTypeN, AIGenTypeN, AIGenTypeN], Attr.Const>;
 688   def : Builtin<name, [AIGenTypeNNoScalar, AIGenTypeNNoScalar, AIGenType1], Attr.Const>;
 689 }
 690 foreach name = ["upsample"] in {
 691   def : Builtin<name, [GenTypeShortVecAndScalar, GenTypeCharVecAndScalar, GenTypeUCharVecAndScalar], Attr.Const>;
 692   def : Builtin<name, [GenTypeUShortVecAndScalar, GenTypeUCharVecAndScalar, GenTypeUCharVecAndScalar], Attr.Const>;
 693   def : Builtin<name, [GenTypeIntVecAndScalar, GenTypeShortVecAndScalar, GenTypeUShortVecAndScalar], Attr.Const>;
 694   def : Builtin<name, [GenTypeUIntVecAndScalar, GenTypeUShortVecAndScalar, GenTypeUShortVecAndScalar], Attr.Const>;
 695   def : Builtin<name, [GenTypeLongVecAndScalar, GenTypeIntVecAndScalar, GenTypeUIntVecAndScalar], Attr.Const>;
 696   def : Builtin<name, [GenTypeULongVecAndScalar, GenTypeUIntVecAndScalar, GenTypeUIntVecAndScalar], Attr.Const>;
 697 }
 698
 699 // --- 3 arguments ---
 700 foreach name = ["clamp"] in {
 701   def : Builtin<name, [AIGenTypeN, AIGenTypeN, AIGenTypeN, AIGenTypeN], Attr.Const>;
 702   def : Builtin<name, [AIGenTypeNNoScalar, AIGenTypeNNoScalar, AIGenType1, AIGenType1], Attr.Const>;
 703 }
 704 foreach name = ["mad_hi", "mad_sat"] in {
 705   def : Builtin<name, [AIGenTypeN, AIGenTypeN, AIGenTypeN, AIGenTypeN], Attr.Const>;
 706 }
 707
 708 // --- Table 11 ---
 709 foreach name = ["mad24"] in {
 710   def : Builtin<name, [GenTypeIntVecAndScalar, GenTypeIntVecAndScalar, GenTypeIntVecAndScalar, GenTypeIntVecAndScalar], Attr.Const>;
 711   def : Builtin<name, [GenTypeUIntVecAndScalar, GenTypeUIntVecAndScalar, GenTypeUIntVecAndScalar, GenTypeUIntVecAndScalar], Attr.Const>;
 712 }
 713 foreach name = ["mul24"] in {
 714   def : Builtin<name, [GenTypeIntVecAndScalar, GenTypeIntVecAndScalar, GenTypeIntVecAndScalar], Attr.Const>;
 715   def : Builtin<name, [GenTypeUIntVecAndScalar, GenTypeUIntVecAndScalar, GenTypeUIntVecAndScalar], Attr.Const>;
 716 }
 717
 718 //--------------------------------------------------------------------
 719 // OpenCL v1.1 s6.11.4, v1.2 s6.12.4, v2.0 s6.13.4 - Common Functions
 720 // OpenCL Extension v2.0 s5.1.3 and s6.1.3 - Common Functions
 721 // --- Table 12 ---
 722 // --- 1 argument ---
 723 foreach name = ["degrees", "radians", "sign"] in {
 724   def : Builtin<name, [FGenTypeN, FGenTypeN], Attr.Const>;
 725 }
 726
 727 // --- 2 arguments ---
 728 foreach name = ["max", "min"] in {
 729   def : Builtin<name, [FGenTypeN, FGenTypeN, FGenTypeN], Attr.Const>;
 730   def : Builtin<name, [GenTypeFloatVecNoScalar, GenTypeFloatVecNoScalar, Float], Attr.Const>;
 731   def : Builtin<name, [GenTypeDoubleVecNoScalar, GenTypeDoubleVecNoScalar, Double], Attr.Const>;
 732   def : Builtin<name, [GenTypeHalfVecNoScalar, GenTypeHalfVecNoScalar, Half], Attr.Const>;
 733 }
 734 foreach name = ["step"] in {
 735   def : Builtin<name, [FGenTypeN, FGenTypeN, FGenTypeN], Attr.Const>;
 736   def : Builtin<name, [GenTypeFloatVecNoScalar, Float, GenTypeFloatVecNoScalar], Attr.Const>;
 737   def : Builtin<name, [GenTypeDoubleVecNoScalar, Double, GenTypeDoubleVecNoScalar], Attr.Const>;
 738   def : Builtin<name, [GenTypeHalfVecNoScalar, Half, GenTypeHalfVecNoScalar], Attr.Const>;
 739 }
 740
 741 // --- 3 arguments ---
 742 foreach name = ["clamp"] in {
 743   def : Builtin<name, [FGenTypeN, FGenTypeN, FGenTypeN, FGenTypeN], Attr.Const>;
 744   def : Builtin<name, [GenTypeFloatVecNoScalar, GenTypeFloatVecNoScalar, Float, Float], Attr.Const>;
 745   def : Builtin<name, [GenTypeDoubleVecNoScalar, GenTypeDoubleVecNoScalar, Double, Double], Attr.Const>;
 746   def : Builtin<name, [GenTypeHalfVecNoScalar, GenTypeHalfVecNoScalar, Half, Half], Attr.Const>;
 747 }
 748 foreach name = ["mix"] in {
 749   def : Builtin<name, [FGenTypeN, FGenTypeN, FGenTypeN, FGenTypeN], Attr.Const>;
 750   def : Builtin<name, [GenTypeFloatVecNoScalar, GenTypeFloatVecNoScalar, GenTypeFloatVecNoScalar, Float], Attr.Const>;
 751   def : Builtin<name, [GenTypeDoubleVecNoScalar, GenTypeDoubleVecNoScalar, GenTypeDoubleVecNoScalar, Double], Attr.Const>;
 752   def : Builtin<name, [GenTypeHalfVecNoScalar, GenTypeHalfVecNoScalar, GenTypeHalfVecNoScalar, Half], Attr.Const>;
 753 }
 754 foreach name = ["smoothstep"] in {
 755   def : Builtin<name, [FGenTypeN, FGenTypeN, FGenTypeN, FGenTypeN], Attr.Const>;
 756   def : Builtin<name, [GenTypeFloatVecNoScalar, Float, Float, GenTypeFloatVecNoScalar], Attr.Const>;
 757   def : Builtin<name, [GenTypeDoubleVecNoScalar, Double, Double, GenTypeDoubleVecNoScalar], Attr.Const>;
 758   def : Builtin<name, [GenTypeHalfVecNoScalar, Half, Half, GenTypeHalfVecNoScalar], Attr.Const>;
 759 }
 760
 761
 762 //--------------------------------------------------------------------
 763 // OpenCL v1.1 s6.11.5, v1.2 s6.12.5, v2.0 s6.13.5 - Geometric Functions
 764 // OpenCL Extension v2.0 s5.1.4 and s6.1.4 - Geometric Functions
 765 // --- Table 13 ---
 766 // --- 1 argument ---
 767 foreach name = ["length"] in {
 768   def : Builtin<name, [Float, GenTypeFloatVec1234], Attr.Const>;
 769   def : Builtin<name, [Double, GenTypeDoubleVec1234], Attr.Const>;
 770   def : Builtin<name, [Half, GenTypeHalfVec1234], Attr.Const>;
 771 }
 772 foreach name = ["normalize"] in {
 773   def : Builtin<name, [GenTypeFloatVec1234, GenTypeFloatVec1234], Attr.Const>;
 774   def : Builtin<name, [GenTypeDoubleVec1234, GenTypeDoubleVec1234], Attr.Const>;
 775   def : Builtin<name, [GenTypeHalfVec1234, GenTypeHalfVec1234], Attr.Const>;
 776 }
 777 foreach name = ["fast_length"] in {
 778   def : Builtin<name, [Float, GenTypeFloatVec1234], Attr.Const>;
 779 }
 780 foreach name = ["fast_normalize"] in {
 781   def : Builtin<name, [GenTypeFloatVec1234, GenTypeFloatVec1234], Attr.Const>;
 782 }
 783
 784 // --- 2 arguments ---
 785 foreach name = ["cross"] in {
 786   foreach VSize = [3, 4] in {
 787     def : Builtin<name, [VectorType<Float, VSize>, VectorType<Float, VSize>, VectorType<Float, VSize>], Attr.Const>;
 788     def : Builtin<name, [VectorType<Double, VSize>, VectorType<Double, VSize>, VectorType<Double, VSize>], Attr.Const>;
 789     def : Builtin<name, [VectorType<Half, VSize>, VectorType<Half, VSize>, VectorType<Half, VSize>], Attr.Const>;
 790   }
 791 }
 792 foreach name = ["dot", "distance"] in {
 793   def : Builtin<name, [Float, GenTypeFloatVec1234, GenTypeFloatVec1234], Attr.Const>;
 794   def : Builtin<name, [Double, GenTypeDoubleVec1234, GenTypeDoubleVec1234], Attr.Const>;
 795   def : Builtin<name, [Half, GenTypeHalfVec1234, GenTypeHalfVec1234], Attr.Const>;
 796 }
 797 foreach name = ["fast_distance"] in {
 798   def : Builtin<name, [Float, GenTypeFloatVec1234, GenTypeFloatVec1234], Attr.Const>;
 799 }
 800
 801
 802 //--------------------------------------------------------------------
 803 // OpenCL v1.1 s6.11.6, v1.2 s6.12.6, v2.0 s6.13.6 - Relational Functions
 804 // OpenCL Extension v2.0 s5.1.5 and s6.1.5 - Relational Functions
 805 // --- Table 14 ---
 806 // --- 1 argument ---
 807 foreach name = ["isfinite", "isinf", "isnan", "isnormal", "signbit"] in {
 808   def : Builtin<name, [GenTypeIntVecAndScalar, GenTypeFloatVecAndScalar], Attr.Const>;
 809   def : Builtin<name, [Int, Double], Attr.Const>;
 810   def : Builtin<name, [GenTypeLongVecNoScalar, GenTypeDoubleVecNoScalar], Attr.Const>;
 811   def : Builtin<name, [Int, Half], Attr.Const>;
 812   def : Builtin<name, [GenTypeShortVecNoScalar, GenTypeHalfVecNoScalar], Attr.Const>;
 813 }
 814 foreach name = ["any", "all"] in {
 815   def : Builtin<name, [Int, SGenTypeN], Attr.Const>;
 816 }
 817
 818 // --- 2 arguments ---
 819 foreach name = ["isequal", "isnotequal", "isgreater", "isgreaterequal",
 820                 "isless", "islessequal", "islessgreater", "isordered",
 821                 "isunordered"] in {
 822   def : Builtin<name, [GenTypeIntVecAndScalar, GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar], Attr.Const>;
 823   def : Builtin<name, [Int, Double, Double], Attr.Const>;
 824   def : Builtin<name, [GenTypeLongVecNoScalar, GenTypeDoubleVecNoScalar, GenTypeDoubleVecNoScalar], Attr.Const>;
 825   def : Builtin<name, [Int, Half, Half], Attr.Const>;
 826   def : Builtin<name, [GenTypeShortVecNoScalar, GenTypeHalfVecNoScalar, GenTypeHalfVecNoScalar], Attr.Const>;
 827 }
 828
 829 // --- 3 arguments ---
 830 foreach name = ["bitselect"] in {
 831   def : Builtin<name, [AGenTypeN, AGenTypeN, AGenTypeN, AGenTypeN], Attr.Const>;
 832 }
 833 foreach name = ["select"] in {
 834   def : Builtin<name, [SGenTypeN, SGenTypeN, SGenTypeN, SGenTypeN], Attr.Const>;
 835   def : Builtin<name, [SGenTypeN, SGenTypeN, SGenTypeN, UGenTypeN], Attr.Const>;
 836   def : Builtin<name, [UGenTypeN, UGenTypeN, UGenTypeN, UGenTypeN], Attr.Const>;
 837   def : Builtin<name, [UGenTypeN, UGenTypeN, UGenTypeN, SGenTypeN], Attr.Const>;
 838   def : Builtin<name, [GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar, GenTypeIntVecAndScalar], Attr.Const>;
 839   def : Builtin<name, [GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar, GenTypeUIntVecAndScalar], Attr.Const>;
 840   def : Builtin<name, [GenTypeDoubleVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeLongVecAndScalar], Attr.Const>;
 841   def : Builtin<name, [GenTypeDoubleVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeULongVecAndScalar], Attr.Const>;
 842   def : Builtin<name, [GenTypeHalfVecAndScalar, GenTypeHalfVecAndScalar, GenTypeHalfVecAndScalar, GenTypeShortVecAndScalar], Attr.Const>;
 843   def : Builtin<name, [GenTypeHalfVecAndScalar, GenTypeHalfVecAndScalar, GenTypeHalfVecAndScalar, GenTypeUShortVecAndScalar], Attr.Const>;
 844 }
 845
 846
 847 //--------------------------------------------------------------------
 848 // OpenCL v1.1 s6.11.7, v1.2 s6.12.7, v2.0 s6.13.7 - Vector Data Load and Store Functions
 849 // OpenCL Extension v1.1 s9.3.6 and s9.6.6, v1.2 s9.5.6, v2.0 s5.1.6 and s6.1.6 - Vector Data Load and Store Functions
 850 // --- Table 15 ---
 851 multiclass VloadVstore<list<AddressSpace> addrspaces, bit defStores> {
 852   foreach AS = addrspaces in {
 853     foreach VSize = [2, 3, 4, 8, 16] in {
 854       foreach name = ["vload" # VSize] in {
 855         def : Builtin<name, [VectorType<Char, VSize>, Size, PointerType<ConstType<Char>, AS>], Attr.Pure>;
 856         def : Builtin<name, [VectorType<UChar, VSize>, Size, PointerType<ConstType<UChar>, AS>], Attr.Pure>;
 857         def : Builtin<name, [VectorType<Short, VSize>, Size, PointerType<ConstType<Short>, AS>], Attr.Pure>;
 858         def : Builtin<name, [VectorType<UShort, VSize>, Size, PointerType<ConstType<UShort>, AS>], Attr.Pure>;
 859         def : Builtin<name, [VectorType<Int, VSize>, Size, PointerType<ConstType<Int>, AS>], Attr.Pure>;
 860         def : Builtin<name, [VectorType<UInt, VSize>, Size, PointerType<ConstType<UInt>, AS>], Attr.Pure>;
 861         def : Builtin<name, [VectorType<Long, VSize>, Size, PointerType<ConstType<Long>, AS>], Attr.Pure>;
 862         def : Builtin<name, [VectorType<ULong, VSize>, Size, PointerType<ConstType<ULong>, AS>], Attr.Pure>;
 863         def : Builtin<name, [VectorType<Float, VSize>, Size, PointerType<ConstType<Float>, AS>], Attr.Pure>;
 864         def : Builtin<name, [VectorType<Double, VSize>, Size, PointerType<ConstType<Double>, AS>], Attr.Pure>;
 865         def : Builtin<name, [VectorType<Half, VSize>, Size, PointerType<ConstType<Half>, AS>], Attr.Pure>;
 866       }
 867       if defStores then {
 868         foreach name = ["vstore" # VSize] in {
 869           def : Builtin<name, [Void, VectorType<Char, VSize>, Size, PointerType<Char, AS>]>;
 870           def : Builtin<name, [Void, VectorType<UChar, VSize>, Size, PointerType<UChar, AS>]>;
 871           def : Builtin<name, [Void, VectorType<Short, VSize>, Size, PointerType<Short, AS>]>;
 872           def : Builtin<name, [Void, VectorType<UShort, VSize>, Size, PointerType<UShort, AS>]>;
 873           def : Builtin<name, [Void, VectorType<Int, VSize>, Size, PointerType<Int, AS>]>;
 874           def : Builtin<name, [Void, VectorType<UInt, VSize>, Size, PointerType<UInt, AS>]>;
 875           def : Builtin<name, [Void, VectorType<Long, VSize>, Size, PointerType<Long, AS>]>;
 876           def : Builtin<name, [Void, VectorType<ULong, VSize>, Size, PointerType<ULong, AS>]>;
 877           def : Builtin<name, [Void, VectorType<Float, VSize>, Size, PointerType<Float, AS>]>;
 878           def : Builtin<name, [Void, VectorType<Double, VSize>, Size, PointerType<Double, AS>]>;
 879           def : Builtin<name, [Void, VectorType<Half, VSize>, Size, PointerType<Half, AS>]>;
 880         }
 881       }
 882     }
 883   }
 884 }
 885
 886 let Extension = FuncExtOpenCLCNamedAddressSpaceBuiltins in {
 887   defm : VloadVstore<[GlobalAS, LocalAS, PrivateAS], 1>;
 888 }
 889 let Extension = FuncExtOpenCLCGenericAddressSpace in {
 890   defm : VloadVstore<[GenericAS], 1>;
 891 }
 892 // vload with constant address space is available regardless of version.
 893 defm : VloadVstore<[ConstantAS], 0>;
 894
 895 multiclass VloadVstoreHalf<list<AddressSpace> addrspaces, bit defStores> {
 896   foreach AS = addrspaces in {
 897     def : Builtin<"vload_half", [Float, Size, !cast<Type>("HalfPtrConst" # AS)], Attr.Pure>;
 898     foreach VSize = [2, 3, 4, 8, 16] in {
 899       foreach name = ["vload_half" # VSize, "vloada_half" # VSize] in {
 900         def : Builtin<name, [VectorType<Float, VSize>, Size, !cast<Type>("HalfPtrConst" # AS)], Attr.Pure>;
 901       }
 902     }
 903     if defStores then {
 904       foreach rnd = ["", "_rte", "_rtz", "_rtp", "_rtn"] in {
 905         foreach name = ["vstore_half" # rnd] in {
 906           def : Builtin<name, [Void, Float, Size, !cast<Type>("HalfPtr" # AS)]>;
 907           def : Builtin<name, [Void, Double, Size, !cast<Type>("HalfPtr" # AS)]>;
 908         }
 909         foreach VSize = [2, 3, 4, 8, 16] in {
 910           foreach name = ["vstore_half" # VSize # rnd, "vstorea_half" # VSize # rnd] in {
 911             def : Builtin<name, [Void, VectorType<Float, VSize>, Size, !cast<Type>("HalfPtr" # AS)]>;
 912             def : Builtin<name, [Void, VectorType<Double, VSize>, Size, !cast<Type>("HalfPtr" # AS)]>;
 913           }
 914         }
 915       }
 916     }
 917   }
 918 }
 919
 920 let Extension = FuncExtOpenCLCNamedAddressSpaceBuiltins in {
 921   defm : VloadVstoreHalf<[GlobalAS, LocalAS, PrivateAS], 1>;
 922 }
 923 let Extension = FuncExtOpenCLCGenericAddressSpace in {
 924   defm : VloadVstoreHalf<[GenericAS], 1>;
 925 }
 926 // vload_half and vloada_half with constant address space are available regardless of version.
 927 defm : VloadVstoreHalf<[ConstantAS], 0>;
 928
 929 // OpenCL v3.0 s6.15.8 - Synchronization Functions.
 930 def : Builtin<"barrier", [Void, MemFenceFlags], Attr.Convergent>;
 931 let MinVersion = CL20 in {
 932   def : Builtin<"work_group_barrier", [Void, MemFenceFlags], Attr.Convergent>;
 933   def : Builtin<"work_group_barrier", [Void, MemFenceFlags, MemoryScope], Attr.Convergent>;
 934 }
 935
 936 // OpenCL v3.0 s6.15.9 - Legacy Explicit Memory Fence Functions.
 937 def : Builtin<"mem_fence", [Void, MemFenceFlags]>;
 938 def : Builtin<"read_mem_fence", [Void, MemFenceFlags]>;
 939 def : Builtin<"write_mem_fence", [Void, MemFenceFlags]>;
 940
 941 // OpenCL v3.0 s6.15.10 - Address Space Qualifier Functions.
 942 // to_global, to_local, to_private are declared in Builtins.def.
 943
 944 let Extension = FuncExtOpenCLCGenericAddressSpace in {
 945   // The OpenCL 3.0 specification defines these with a "gentype" argument indicating any builtin
 946   // type or user-defined type, which cannot be represented currently.  Hence we slightly diverge
 947   // by providing only the following overloads with a void pointer.
 948   def : Builtin<"get_fence", [MemFenceFlags, PointerType<Void, GenericAS>]>;
 949   def : Builtin<"get_fence", [MemFenceFlags, PointerType<ConstType<Void>, GenericAS>]>;
 950 }
 951
 952 //--------------------------------------------------------------------
 953 // OpenCL v1.1 s6.11.10, v1.2 s6.12.10, v2.0 s6.13.10: Async Copies from Global to Local Memory, Local to Global Memory, and Prefetch
 954 // OpenCL Extension v2.0 s5.1.7 and s6.1.7: Async Copies from Global to Local Memory, Local to Global Memory, and Prefetch
 955 // --- Table 18 ---
 956 foreach name = ["async_work_group_copy"] in {
 957   def : Builtin<name, [Event, PointerType<AGenTypeN, LocalAS>, PointerType<ConstType<AGenTypeN>, GlobalAS>, Size, Event]>;
 958   def : Builtin<name, [Event, PointerType<AGenTypeN, GlobalAS>, PointerType<ConstType<AGenTypeN>, LocalAS>, Size, Event]>;
 959 }
 960 foreach name = ["async_work_group_strided_copy"] in {
 961   def : Builtin<name, [Event, PointerType<AGenTypeN, LocalAS>, PointerType<ConstType<AGenTypeN>, GlobalAS>, Size, Size, Event]>;
 962   def : Builtin<name, [Event, PointerType<AGenTypeN, GlobalAS>, PointerType<ConstType<AGenTypeN>, LocalAS>, Size, Size, Event]>;
 963 }
 964 foreach name = ["wait_group_events"] in {
 965   def : Builtin<name, [Void, Int, PointerType<Event, GenericAS>]>;
 966 }
 967 foreach name = ["prefetch"] in {
 968   def : Builtin<name, [Void, PointerType<ConstType<AGenTypeN>, GlobalAS>, Size]>;
 969 }
 970
 971 //--------------------------------------------------------------------
 972 // OpenCL v2.0 s6.13.11 - Atomics Functions.
 973 // Functions that use memory_order and cl_mem_fence_flags enums are not
 974 // declared here as the TableGen backend does not handle enums.
 975
 976 // OpenCL v1.0 s9.5, s9.6, s9.7 - Atomic Functions for 32-bit integers
 977 // --- Table 9.1 ---
 978 let Extension = FuncExtKhrGlobalInt32BaseAtomics in {
 979   foreach Type = [Int, UInt] in {
 980     foreach name = ["atom_add", "atom_sub", "atom_xchg"] in {
 981       def : Builtin<name, [Type, PointerType<VolatileType<Type>, GlobalAS>, Type]>;
 982     }
 983     foreach name = ["atom_inc", "atom_dec"] in {
 984       def : Builtin<name, [Type, PointerType<VolatileType<Type>, GlobalAS>]>;
 985     }
 986     foreach name = ["atom_cmpxchg"] in {
 987       def : Builtin<name, [Type, PointerType<VolatileType<Type>, GlobalAS>, Type, Type]>;
 988     }
 989   }
 990 }
 991 // --- Table 9.3 ---
 992 let Extension = FuncExtKhrLocalInt32BaseAtomics in {
 993   foreach Type = [Int, UInt] in {
 994     foreach name = ["atom_add", "atom_sub", "atom_xchg"] in {
 995       def : Builtin<name, [Type, PointerType<VolatileType<Type>, LocalAS>, Type]>;
 996     }
 997     foreach name = ["atom_inc", "atom_dec"] in {
 998       def : Builtin<name, [Type, PointerType<VolatileType<Type>, LocalAS>]>;
 999     }
1000     foreach name = ["atom_cmpxchg"] in {
1001       def : Builtin<name, [Type, PointerType<VolatileType<Type>, LocalAS>, Type, Type]>;
1002     }
1003   }
1004 }
1005 // --- Table 9.5 ---
1006 let Extension = FuncExtKhrInt64BaseAtomics in {
1007   foreach AS = [GlobalAS, LocalAS] in {
1008     foreach Type = [Long, ULong] in {
1009       foreach name = ["atom_add", "atom_sub", "atom_xchg"] in {
1010         def : Builtin<name, [Type, PointerType<VolatileType<Type>, AS>, Type]>;
1011       }
1012       foreach name = ["atom_inc", "atom_dec"] in {
1013         def : Builtin<name, [Type, PointerType<VolatileType<Type>, AS>]>;
1014       }
1015       foreach name = ["atom_cmpxchg"] in {
1016         def : Builtin<name, [Type, PointerType<VolatileType<Type>, AS>, Type, Type]>;
1017       }
1018     }
1019   }
1020 }
1021 // --- Table 9.2 ---
1022 let Extension = FuncExtKhrGlobalInt32ExtendedAtomics in {
1023   foreach Type = [Int, UInt] in {
1024     foreach name = ["atom_min", "atom_max", "atom_and",
1025                     "atom_or", "atom_xor"] in {
1026       def : Builtin<name, [Type, PointerType<VolatileType<Type>, GlobalAS>, Type]>;
1027     }
1028   }
1029 }
1030 // --- Table 9.4 ---
1031 let Extension = FuncExtKhrLocalInt32ExtendedAtomics in {
1032   foreach Type = [Int, UInt] in {
1033     foreach name = ["atom_min", "atom_max", "atom_and",
1034                     "atom_or", "atom_xor"] in {
1035       def : Builtin<name, [Type, PointerType<VolatileType<Type>, LocalAS>, Type]>;
1036     }
1037   }
1038 }
1039 // --- Table 9.6 ---
1040 let Extension = FuncExtKhrInt64ExtendedAtomics in {
1041   foreach AS = [GlobalAS, LocalAS] in {
1042     foreach Type = [Long, ULong] in {
1043       foreach name = ["atom_min", "atom_max", "atom_and",
1044                       "atom_or", "atom_xor"] in {
1045         def : Builtin<name, [Type, PointerType<VolatileType<Type>, AS>, Type]>;
1046       }
1047     }
1048   }
1049 }
1050 // OpenCL v1.1 s6.11.1, v1.2 s6.12.11 - Atomic Functions
1051 foreach AS = [GlobalAS, LocalAS] in {
1052   def : Builtin<"atomic_xchg", [Float, PointerType<VolatileType<Float>, AS>, Float]>;
1053   foreach Type = [Int, UInt] in {
1054     foreach name = ["atomic_add", "atomic_sub", "atomic_xchg",
1055                     "atomic_min", "atomic_max", "atomic_and",
1056                     "atomic_or", "atomic_xor"] in {
1057       def : Builtin<name, [Type, PointerType<VolatileType<Type>, AS>, Type]>;
1058     }
1059     foreach name = ["atomic_inc", "atomic_dec"] in {
1060       def : Builtin<name, [Type, PointerType<VolatileType<Type>, AS>]>;
1061     }
1062     foreach name = ["atomic_cmpxchg"] in {
1063       def : Builtin<name, [Type, PointerType<VolatileType<Type>, AS>, Type, Type]>;
1064     }
1065   }
1066 }
1067
1068 let Extension = FuncExtOpenCLCxx in {
1069   foreach Type = [Int, UInt] in {
1070     foreach name = ["atomic_add", "atomic_sub", "atomic_xchg",
1071                     "atomic_min", "atomic_max", "atomic_and",
1072                     "atomic_or", "atomic_xor"] in {
1073       def : Builtin<name, [Type, PointerType<VolatileType<Type>, GenericAS>, Type]>;
1074     }
1075     foreach name = ["atomic_inc", "atomic_dec"] in {
1076       def : Builtin<name, [Type, PointerType<VolatileType<Type>, GenericAS>]>;
1077     }
1078     foreach name = ["atomic_cmpxchg"] in {
1079       def : Builtin<name, [Type, PointerType<VolatileType<Type>, GenericAS>, Type, Type]>;
1080     }
1081   }
1082 }
1083
1084 // OpenCL v2.0 s6.13.11 - Atomic Functions.
1085
1086 // An atomic builtin with 2 additional _explicit variants.
1087 multiclass BuiltinAtomicExplicit<string Name, list<Type> Types, FunctionExtension BaseExt> {
1088   // Without explicit MemoryOrder or MemoryScope.
1089   let Extension = concatExtension<BaseExt, "__opencl_c_atomic_order_seq_cst __opencl_c_atomic_scope_device">.ret in {
1090     def : Builtin<Name, Types>;
1091   }
1092
1093   // With an explicit MemoryOrder argument.
1094   let Extension = concatExtension<BaseExt, "__opencl_c_atomic_scope_device">.ret in {
1095     def : Builtin<Name # "_explicit", !listconcat(Types, [MemoryOrder])>;
1096   }
1097
1098   // With explicit MemoryOrder and MemoryScope arguments.
1099   let Extension = BaseExt in {
1100     def : Builtin<Name # "_explicit", !listconcat(Types, [MemoryOrder, MemoryScope])>;
1101   }
1102 }
1103
1104 // OpenCL 2.0 atomic functions that have a pointer argument in a given address space.
1105 multiclass OpenCL2Atomics<AddressSpace addrspace, FunctionExtension BaseExt> {
1106   foreach TypePair = [[AtomicInt, Int], [AtomicUInt, UInt],
1107                       [AtomicLong, Long], [AtomicULong, ULong],
1108                       [AtomicFloat, Float], [AtomicDouble, Double]] in {
1109     let Extension = BaseExt in {
1110       def : Builtin<"atomic_init",
1111           [Void, PointerType<VolatileType<TypePair[0]>, addrspace>, TypePair[1]]>;
1112     }
1113     defm : BuiltinAtomicExplicit<"atomic_store",
1114         [Void, PointerType<VolatileType<TypePair[0]>, addrspace>, TypePair[1]], BaseExt>;
1115     defm : BuiltinAtomicExplicit<"atomic_load",
1116         [TypePair[1], PointerType<VolatileType<TypePair[0]>, addrspace>], BaseExt>;
1117     defm : BuiltinAtomicExplicit<"atomic_exchange",
1118         [TypePair[1], PointerType<VolatileType<TypePair[0]>, addrspace>, TypePair[1]], BaseExt>;
1119     foreach Variant = ["weak", "strong"] in {
1120       foreach exp_ptr_addrspace = !cond(
1121             !eq(BaseExt, FuncExtOpenCLCGenericAddressSpace): [GenericAS],
1122             !eq(BaseExt, FuncExtOpenCLCNamedAddressSpaceBuiltins): [GlobalAS, LocalAS, PrivateAS])
1123           in {
1124         let Extension = concatExtension<BaseExt, "__opencl_c_atomic_order_seq_cst __opencl_c_atomic_scope_device">.ret in {
1125           def : Builtin<"atomic_compare_exchange_" # Variant,
1126               [Bool, PointerType<VolatileType<TypePair[0]>, addrspace>,
1127                PointerType<TypePair[1], exp_ptr_addrspace>, TypePair[1]]>;
1128         }
1129         let Extension = concatExtension<BaseExt, "__opencl_c_atomic_scope_device">.ret in {
1130           def : Builtin<"atomic_compare_exchange_" # Variant # "_explicit",
1131               [Bool, PointerType<VolatileType<TypePair[0]>, addrspace>,
1132                PointerType<TypePair[1], exp_ptr_addrspace>, TypePair[1], MemoryOrder, MemoryOrder]>;
1133         }
1134         let Extension = BaseExt in {
1135           def : Builtin<"atomic_compare_exchange_" # Variant # "_explicit",
1136               [Bool, PointerType<VolatileType<TypePair[0]>, addrspace>,
1137                PointerType<TypePair[1], exp_ptr_addrspace>, TypePair[1], MemoryOrder, MemoryOrder, MemoryScope]>;
1138         }
1139       }
1140     }
1141   }
1142
1143   foreach TypePair = [[AtomicInt, Int, Int], [AtomicUInt, UInt, UInt],
1144                       [AtomicLong, Long, Long], [AtomicULong, ULong, ULong],
1145                       [AtomicUIntPtr, UIntPtr, PtrDiff]] in {
1146     foreach ModOp = ["add", "sub"] in {
1147       defm : BuiltinAtomicExplicit<"atomic_fetch_" # ModOp,
1148           [TypePair[1], PointerType<VolatileType<TypePair[0]>, addrspace>, TypePair[2]], BaseExt>;
1149     }
1150   }
1151   foreach TypePair = [[AtomicInt, Int, Int], [AtomicUInt, UInt, UInt],
1152                       [AtomicLong, Long, Long], [AtomicULong, ULong, ULong]] in {
1153     foreach ModOp = ["or", "xor", "and", "min", "max"] in {
1154       defm : BuiltinAtomicExplicit<"atomic_fetch_" # ModOp,
1155           [TypePair[1], PointerType<VolatileType<TypePair[0]>, addrspace>, TypePair[2]], BaseExt>;
1156     }
1157   }
1158
1159   defm : BuiltinAtomicExplicit<"atomic_flag_clear",
1160       [Void, PointerType<VolatileType<AtomicFlag>, addrspace>], BaseExt>;
1161
1162   defm : BuiltinAtomicExplicit<"atomic_flag_test_and_set",
1163       [Bool, PointerType<VolatileType<AtomicFlag>, addrspace>], BaseExt>;
1164 }
1165
1166 let MinVersion = CL20 in {
1167   def : Builtin<"atomic_work_item_fence", [Void, MemFenceFlags, MemoryOrder, MemoryScope]>;
1168
1169   defm : OpenCL2Atomics<GenericAS, FuncExtOpenCLCGenericAddressSpace>;
1170   defm : OpenCL2Atomics<GlobalAS, FuncExtOpenCLCNamedAddressSpaceBuiltins>;
1171   defm : OpenCL2Atomics<LocalAS, FuncExtOpenCLCNamedAddressSpaceBuiltins>;
1172 }
1173
1174 // The functionality added by cl_ext_float_atomics extension
1175 let MinVersion = CL20 in {
1176   foreach addrspace = [GlobalAS, LocalAS, GenericAS] in {
1177     defvar extension_fp16 = !cast<FunctionExtension>("FuncExtFloatAtomicsFp16" # addrspace # "LoadStore");
1178
1179     defm : BuiltinAtomicExplicit<"atomic_store",
1180         [Void, PointerType<VolatileType<AtomicHalf>, addrspace>, AtomicHalf], extension_fp16>;
1181     defm : BuiltinAtomicExplicit<"atomic_load",
1182         [Half, PointerType<VolatileType<AtomicHalf>, addrspace>], extension_fp16>;
1183     defm : BuiltinAtomicExplicit<"atomic_exchange",
1184         [Half, PointerType<VolatileType<AtomicHalf>, addrspace>, Half], extension_fp16>;
1185
1186     foreach ModOp = ["add", "sub"] in {
1187       defvar extension_fp16 = !cast<FunctionExtension>("FuncExtFloatAtomicsFp16" # addrspace # "Add");
1188       defvar extension_fp32 = !cast<FunctionExtension>("FuncExtFloatAtomicsFp32" # addrspace # "Add");
1189       defvar extension_fp64 = !cast<FunctionExtension>("FuncExtFloatAtomicsFp64" # addrspace # "Add");
1190
1191       defm : BuiltinAtomicExplicit<"atomic_fetch_" # ModOp,
1192           [Half, PointerType<VolatileType<AtomicHalf>, addrspace>, Half], extension_fp16>;
1193       defm : BuiltinAtomicExplicit<"atomic_fetch_" # ModOp,
1194           [Float, PointerType<VolatileType<AtomicFloat>, addrspace>, Float], extension_fp32>;
1195       defm : BuiltinAtomicExplicit<"atomic_fetch_" # ModOp,
1196           [Double, PointerType<VolatileType<AtomicDouble>, addrspace>, Double], extension_fp64>;
1197     }
1198
1199     foreach ModOp = ["min", "max"] in {
1200       defvar extension_fp16 = !cast<FunctionExtension>("FuncExtFloatAtomicsFp16" # addrspace # "MinMax");
1201       defvar extension_fp32 = !cast<FunctionExtension>("FuncExtFloatAtomicsFp32" # addrspace # "MinMax");
1202       defvar extension_fp64 = !cast<FunctionExtension>("FuncExtFloatAtomicsFp64" # addrspace # "MinMax");
1203
1204       defm : BuiltinAtomicExplicit<"atomic_fetch_" # ModOp,
1205           [Half, PointerType<VolatileType<AtomicHalf>, addrspace>, Half], extension_fp16>;
1206       defm : BuiltinAtomicExplicit<"atomic_fetch_" # ModOp,
1207           [Float, PointerType<VolatileType<AtomicFloat>, addrspace>, Float], extension_fp32>;
1208       defm : BuiltinAtomicExplicit<"atomic_fetch_" # ModOp,
1209           [Double, PointerType<VolatileType<AtomicDouble>, addrspace>, Double], extension_fp64>;
1210     }
1211   }
1212 }
1213
1214 //--------------------------------------------------------------------
1215 // OpenCL v1.1 s6.11.12, v1.2 s6.12.12, v2.0 s6.13.12 - Miscellaneous Vector Functions
1216 // --- Table 19 ---
1217 foreach VSize1 = [2, 4, 8, 16] in {
1218   foreach VSize2 = [2, 4, 8, 16] in {
1219     foreach VecAndMaskType = [[Char, UChar], [UChar, UChar],
1220                               [Short, UShort], [UShort, UShort],
1221                               [Int, UInt], [UInt, UInt],
1222                               [Long, ULong], [ULong, ULong],
1223                               [Float, UInt], [Double, ULong], [Half, UShort]] in {
1224       def : Builtin<"shuffle", [VectorType<VecAndMaskType[0], VSize1>,
1225                                 VectorType<VecAndMaskType[0], VSize2>,
1226                                 VectorType<VecAndMaskType[1], VSize1>],
1227                                Attr.Const>;
1228     }
1229   }
1230 }
1231 foreach VSize1 = [2, 4, 8, 16] in {
1232   foreach VSize2 = [2, 4, 8, 16] in {
1233     foreach VecAndMaskType = [[Char, UChar], [UChar, UChar],
1234                               [Short, UShort], [UShort, UShort],
1235                               [Int, UInt], [UInt, UInt],
1236                               [Long, ULong], [ULong, ULong],
1237                               [Float, UInt], [Double, ULong], [Half, UShort]] in {
1238       def : Builtin<"shuffle2", [VectorType<VecAndMaskType[0], VSize1>,
1239                                  VectorType<VecAndMaskType[0], VSize2>,
1240                                  VectorType<VecAndMaskType[0], VSize2>,
1241                                  VectorType<VecAndMaskType[1], VSize1>],
1242                                 Attr.Const>;
1243     }
1244   }
1245 }
1246
1247 //--------------------------------------------------------------------
1248 // OpenCL v1.1 s6.11.3, v1.2 s6.12.14, v2.0 s6.13.14: Image Read and Write Functions
1249 // OpenCL Extension v2.0 s5.1.8 and s6.1.8: Image Read and Write Functions
1250 // --- Table 22: Image Read Functions with Samplers ---
1251 foreach imgTy = [Image1d] in {
1252   foreach coordTy = [Int, Float] in {
1253     def : Builtin<"read_imagef", [VectorType<Float, 4>, ImageType<imgTy, "RO">, Sampler, coordTy], Attr.Pure>;
1254     def : Builtin<"read_imagei", [VectorType<Int, 4>, ImageType<imgTy, "RO">, Sampler, coordTy], Attr.Pure>;
1255     def : Builtin<"read_imageui", [VectorType<UInt, 4>, ImageType<imgTy, "RO">, Sampler, coordTy], Attr.Pure>;
1256   }
1257 }
1258 foreach imgTy = [Image2d, Image1dArray] in {
1259   foreach coordTy = [Int, Float] in {
1260     def : Builtin<"read_imagef", [VectorType<Float, 4>, ImageType<imgTy, "RO">, Sampler, VectorType<coordTy, 2>], Attr.Pure>;
1261     def : Builtin<"read_imagei", [VectorType<Int, 4>, ImageType<imgTy, "RO">, Sampler, VectorType<coordTy, 2>], Attr.Pure>;
1262     def : Builtin<"read_imageui", [VectorType<UInt, 4>, ImageType<imgTy, "RO">, Sampler, VectorType<coordTy, 2>], Attr.Pure>;
1263   }
1264 }
1265 foreach imgTy = [Image3d, Image2dArray] in {
1266   foreach coordTy = [Int, Float] in {
1267     def : Builtin<"read_imagef", [VectorType<Float, 4>, ImageType<imgTy, "RO">, Sampler, VectorType<coordTy, 4>], Attr.Pure>;
1268     def : Builtin<"read_imagei", [VectorType<Int, 4>, ImageType<imgTy, "RO">, Sampler, VectorType<coordTy, 4>], Attr.Pure>;
1269     def : Builtin<"read_imageui", [VectorType<UInt, 4>, ImageType<imgTy, "RO">, Sampler, VectorType<coordTy, 4>], Attr.Pure>;
1270   }
1271 }
1272 foreach coordTy = [Int, Float] in {
1273   def : Builtin<"read_imagef", [Float, ImageType<Image2dDepth, "RO">, Sampler, VectorType<coordTy, 2>], Attr.Pure>;
1274   def : Builtin<"read_imagef", [Float, ImageType<Image2dArrayDepth, "RO">, Sampler, VectorType<coordTy, 4>], Attr.Pure>;
1275 }
1276
1277 // --- Table 23: Sampler-less Read Functions ---
1278 multiclass ImageReadSamplerless<string aQual> {
1279   foreach imgTy = [Image2d, Image1dArray] in {
1280     def : Builtin<"read_imagef", [VectorType<Float, 4>, ImageType<imgTy, aQual>, VectorType<Int, 2>], Attr.Pure>;
1281     def : Builtin<"read_imagei", [VectorType<Int, 4>, ImageType<imgTy, aQual>, VectorType<Int, 2>], Attr.Pure>;
1282     def : Builtin<"read_imageui", [VectorType<UInt, 4>, ImageType<imgTy, aQual>, VectorType<Int, 2>], Attr.Pure>;
1283   }
1284   foreach imgTy = [Image3d, Image2dArray] in {
1285     def : Builtin<"read_imagef", [VectorType<Float, 4>, ImageType<imgTy, aQual>, VectorType<Int, 4>], Attr.Pure>;
1286     def : Builtin<"read_imagei", [VectorType<Int, 4>, ImageType<imgTy, aQual>, VectorType<Int, 4>], Attr.Pure>;
1287     def : Builtin<"read_imageui", [VectorType<UInt, 4>, ImageType<imgTy, aQual>, VectorType<Int, 4>], Attr.Pure>;
1288   }
1289   foreach imgTy = [Image1d, Image1dBuffer] in {
1290     def : Builtin<"read_imagef", [VectorType<Float, 4>, ImageType<imgTy, aQual>, Int], Attr.Pure>;
1291     def : Builtin<"read_imagei", [VectorType<Int, 4>, ImageType<imgTy, aQual>, Int], Attr.Pure>;
1292     def : Builtin<"read_imageui", [VectorType<UInt, 4>, ImageType<imgTy, aQual>, Int], Attr.Pure>;
1293   }
1294   def : Builtin<"read_imagef", [Float, ImageType<Image2dDepth, aQual>, VectorType<Int, 2>], Attr.Pure>;
1295   def : Builtin<"read_imagef", [Float, ImageType<Image2dArrayDepth, aQual>, VectorType<Int, 4>], Attr.Pure>;
1296 }
1297
1298 let MinVersion = CL12 in {
1299   defm : ImageReadSamplerless<"RO">;
1300   let Extension = FuncExtOpenCLCReadWriteImages in {
1301     defm : ImageReadSamplerless<"RW">;
1302   }
1303 }
1304
1305 // --- Table 24: Image Write Functions ---
1306 multiclass ImageWrite<string aQual> {
1307   foreach imgTy = [Image2d] in {
1308     def : Builtin<"write_imagef", [Void, ImageType<imgTy, aQual>, VectorType<Int, 2>, VectorType<Float, 4>]>;
1309     def : Builtin<"write_imagei", [Void, ImageType<imgTy, aQual>, VectorType<Int, 2>, VectorType<Int, 4>]>;
1310     def : Builtin<"write_imageui", [Void, ImageType<imgTy, aQual>, VectorType<Int, 2>, VectorType<UInt, 4>]>;
1311   }
1312   foreach imgTy = [Image2dArray] in {
1313     def : Builtin<"write_imagef", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, VectorType<Float, 4>]>;
1314     def : Builtin<"write_imagei", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, VectorType<Int, 4>]>;
1315     def : Builtin<"write_imageui", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, VectorType<UInt, 4>]>;
1316   }
1317   foreach imgTy = [Image1d, Image1dBuffer] in {
1318     def : Builtin<"write_imagef", [Void, ImageType<imgTy, aQual>, Int, VectorType<Float, 4>]>;
1319     def : Builtin<"write_imagei", [Void, ImageType<imgTy, aQual>, Int, VectorType<Int, 4>]>;
1320     def : Builtin<"write_imageui", [Void, ImageType<imgTy, aQual>, Int, VectorType<UInt, 4>]>;
1321   }
1322   foreach imgTy = [Image1dArray] in {
1323     def : Builtin<"write_imagef", [Void, ImageType<imgTy, aQual>, VectorType<Int, 2>, VectorType<Float, 4>]>;
1324     def : Builtin<"write_imagei", [Void, ImageType<imgTy, aQual>, VectorType<Int, 2>, VectorType<Int, 4>]>;
1325     def : Builtin<"write_imageui", [Void, ImageType<imgTy, aQual>, VectorType<Int, 2>, VectorType<UInt, 4>]>;
1326   }
1327   foreach imgTy = [Image3d] in {
1328     def : Builtin<"write_imagef", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, VectorType<Float, 4>]>;
1329     def : Builtin<"write_imagei", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, VectorType<Int, 4>]>;
1330     def : Builtin<"write_imageui", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, VectorType<UInt, 4>]>;
1331   }
1332   def : Builtin<"write_imagef", [Void, ImageType<Image2dDepth, aQual>, VectorType<Int, 2>, Float]>;
1333   def : Builtin<"write_imagef", [Void, ImageType<Image2dArrayDepth, aQual>, VectorType<Int, 4>, Float]>;
1334 }
1335
1336 defm : ImageWrite<"WO">;
1337 let Extension = FuncExtOpenCLCReadWriteImages in {
1338   defm : ImageWrite<"RW">;
1339 }
1340
1341 // --- Table 25: Image Query Functions ---
1342 multiclass ImageQuery<string aQual> {
1343   foreach imgTy = [Image1d, Image1dBuffer, Image2d, Image3d,
1344                    Image1dArray, Image2dArray, Image2dDepth,
1345                    Image2dArrayDepth] in {
1346     foreach name = ["get_image_width", "get_image_channel_data_type",
1347                     "get_image_channel_order"] in {
1348       def : Builtin<name, [Int, ImageType<imgTy, aQual>], Attr.Const>;
1349     }
1350   }
1351   foreach imgTy = [Image2d, Image3d, Image2dArray, Image2dDepth,
1352                    Image2dArrayDepth] in {
1353     def : Builtin<"get_image_height", [Int, ImageType<imgTy, aQual>], Attr.Const>;
1354   }
1355   def : Builtin<"get_image_depth", [Int, ImageType<Image3d, aQual>], Attr.Const>;
1356   foreach imgTy = [Image2d, Image2dArray, Image2dDepth,
1357                    Image2dArrayDepth] in {
1358     def : Builtin<"get_image_dim", [VectorType<Int, 2>, ImageType<imgTy, aQual>], Attr.Const>;
1359   }
1360   def : Builtin<"get_image_dim", [VectorType<Int, 4>, ImageType<Image3d, aQual>], Attr.Const>;
1361   foreach imgTy = [Image1dArray, Image2dArray, Image2dArrayDepth] in {
1362     def : Builtin<"get_image_array_size", [Size, ImageType<imgTy, aQual>], Attr.Const>;
1363   }
1364 }
1365
1366 defm : ImageQuery<"RO">;
1367 defm : ImageQuery<"WO">;
1368 let Extension = FuncExtOpenCLCReadWriteImages in {
1369   defm : ImageQuery<"RW">;
1370 }
1371
1372 // OpenCL extension v2.0 s5.1.9: Built-in Image Read Functions
1373 // --- Table 8 ---
1374 foreach aQual = ["RO"] in {
1375   foreach name = ["read_imageh"] in {
1376     foreach coordTy = [Int, Float] in {
1377       foreach imgTy = [Image2d, Image1dArray] in {
1378         def : Builtin<name, [VectorType<Half, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<coordTy, 2>], Attr.Pure>;
1379       }
1380       foreach imgTy = [Image3d, Image2dArray] in {
1381         def : Builtin<name, [VectorType<Half, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<coordTy, 4>], Attr.Pure>;
1382       }
1383       foreach imgTy = [Image1d] in {
1384         def : Builtin<name, [VectorType<Half, 4>, ImageType<imgTy, aQual>, Sampler, coordTy], Attr.Pure>;
1385       }
1386     }
1387   }
1388 }
1389 // OpenCL extension v2.0 s5.1.10: Built-in Image Sampler-less Read Functions
1390 // --- Table 9 ---
1391 let MinVersion = CL12 in {
1392   multiclass ImageReadHalf<string aQual> {
1393     foreach name = ["read_imageh"] in {
1394       foreach imgTy = [Image2d, Image1dArray] in {
1395         def : Builtin<name, [VectorType<Half, 4>, ImageType<imgTy, aQual>, VectorType<Int, 2>], Attr.Pure>;
1396       }
1397       foreach imgTy = [Image3d, Image2dArray] in {
1398         def : Builtin<name, [VectorType<Half, 4>, ImageType<imgTy, aQual>, VectorType<Int, 4>], Attr.Pure>;
1399       }
1400       foreach imgTy = [Image1d, Image1dBuffer] in {
1401         def : Builtin<name, [VectorType<Half, 4>, ImageType<imgTy, aQual>, Int], Attr.Pure>;
1402       }
1403     }
1404   }
1405   defm : ImageReadHalf<"RO">;
1406   let Extension = FuncExtOpenCLCReadWriteImages in {
1407     defm : ImageReadHalf<"RW">;
1408   }
1409 }
1410 // OpenCL extension v2.0 s5.1.11: Built-in Image Write Functions
1411 // --- Table 10 ---
1412 multiclass ImageWriteHalf<string aQual> {
1413   foreach name = ["write_imageh"] in {
1414     def : Builtin<name, [Void, ImageType<Image2d, aQual>, VectorType<Int, 2>, VectorType<Half, 4>]>;
1415     def : Builtin<name, [Void, ImageType<Image2dArray, aQual>, VectorType<Int, 4>, VectorType<Half, 4>]>;
1416     def : Builtin<name, [Void, ImageType<Image1d, aQual>, Int, VectorType<Half, 4>]>;
1417     def : Builtin<name, [Void, ImageType<Image1dBuffer, aQual>, Int, VectorType<Half, 4>]>;
1418     def : Builtin<name, [Void, ImageType<Image1dArray, aQual>, VectorType<Int, 2>, VectorType<Half, 4>]>;
1419     def : Builtin<name, [Void, ImageType<Image3d, aQual>, VectorType<Int, 4>, VectorType<Half, 4>]>;
1420   }
1421 }
1422
1423 defm : ImageWriteHalf<"WO">;
1424 let Extension = FuncExtOpenCLCReadWriteImages in {
1425   defm : ImageWriteHalf<"RW">;
1426 }
1427
1428
1429
1430 //--------------------------------------------------------------------
1431 // OpenCL v2.0 s6.13.15 - Work-group Functions
1432 // --- Table 26 ---
1433 let Extension = FuncExtOpenCLCWGCollectiveFunctions in {
1434   foreach name = ["work_group_all", "work_group_any"] in {
1435     def : Builtin<name, [Int, Int], Attr.Convergent>;
1436   }
1437   foreach name = ["work_group_broadcast"] in {
1438     def : Builtin<name, [IntLongFloatGenType1, IntLongFloatGenType1, Size], Attr.Convergent>;
1439     def : Builtin<name, [IntLongFloatGenType1, IntLongFloatGenType1, Size, Size], Attr.Convergent>;
1440     def : Builtin<name, [IntLongFloatGenType1, IntLongFloatGenType1, Size, Size, Size], Attr.Convergent>;
1441   }
1442   foreach op = ["add", "min", "max"] in {
1443     foreach name = ["work_group_reduce_", "work_group_scan_exclusive_",
1444                     "work_group_scan_inclusive_"] in {
1445       def : Builtin<name # op, [IntLongFloatGenType1, IntLongFloatGenType1], Attr.Convergent>;
1446     }
1447   }
1448 }
1449
1450
1451 //--------------------------------------------------------------------
1452 // OpenCL2.0 : 6.13.16 : Pipe Functions
1453 // --- Table 27 ---
1454 // Defined in Builtins.def
1455
1456 // --- Table 28 ---
1457 // Builtins taking pipe arguments are defined in Builtins.def
1458 let Extension = FuncExtOpenCLCPipes in {
1459   def : Builtin<"is_valid_reserve_id", [Bool, ReserveId]>;
1460 }
1461
1462 // --- Table 29 ---
1463 // Defined in Builtins.def
1464
1465
1466 //--------------------------------------------------------------------
1467 // OpenCL2.0 : 6.13.17 : Enqueuing Kernels
1468 // --- Table 30 ---
1469 // Defined in Builtins.def
1470
1471 // --- Table 32 ---
1472 // Defined in Builtins.def
1473
1474 // --- Table 33 ---
1475 let Extension = FuncExtOpenCLCDeviceEnqueue in {
1476   def : Builtin<"enqueue_marker",
1477       [Int, Queue, UInt, PointerType<ConstType<ClkEvent>, GenericAS>, PointerType<ClkEvent, GenericAS>]>;
1478
1479   // --- Table 34 ---
1480   def : Builtin<"retain_event", [Void, ClkEvent]>;
1481   def : Builtin<"release_event", [Void, ClkEvent]>;
1482   def : Builtin<"create_user_event", [ClkEvent]>;
1483   def : Builtin<"is_valid_event", [Bool, ClkEvent]>;
1484   def : Builtin<"set_user_event_status", [Void, ClkEvent, Int]>;
1485   def : Builtin<"capture_event_profiling_info",
1486       [Void, ClkEvent, ClkProfilingInfo, PointerType<Void, GlobalAS>]>;
1487
1488   // --- Table 35 ---
1489   def : Builtin<"get_default_queue", [Queue]>;
1490
1491   def : Builtin<"ndrange_1D", [NDRange, Size]>;
1492   def : Builtin<"ndrange_1D", [NDRange, Size, Size]>;
1493   def : Builtin<"ndrange_1D", [NDRange, Size, Size, Size]>;
1494   def : Builtin<"ndrange_2D", [NDRange, PointerType<ConstType<Size>, PrivateAS>]>;
1495   def : Builtin<"ndrange_2D", [NDRange, PointerType<ConstType<Size>, PrivateAS>,
1496                                         PointerType<ConstType<Size>, PrivateAS>]>;
1497   def : Builtin<"ndrange_2D", [NDRange, PointerType<ConstType<Size>, PrivateAS>,
1498                                         PointerType<ConstType<Size>, PrivateAS>,
1499                                         PointerType<ConstType<Size>, PrivateAS>]>;
1500   def : Builtin<"ndrange_3D", [NDRange, PointerType<ConstType<Size>, PrivateAS>]>;
1501   def : Builtin<"ndrange_3D", [NDRange, PointerType<ConstType<Size>, PrivateAS>,
1502                                         PointerType<ConstType<Size>, PrivateAS>]>;
1503   def : Builtin<"ndrange_3D", [NDRange, PointerType<ConstType<Size>, PrivateAS>,
1504                                         PointerType<ConstType<Size>, PrivateAS>,
1505                                         PointerType<ConstType<Size>, PrivateAS>]>;
1506 }
1507
1508
1509 //--------------------------------------------------------------------
1510 // End of the builtin functions defined in the OpenCL C specification.
1511 // Builtin functions defined in the OpenCL C Extension are below.
1512 //--------------------------------------------------------------------
1513
1514
1515 // OpenCL Extension v2.0 s9.18 - Mipmaps
1516 let Extension = FuncExtKhrMipmapImage in {
1517   // Added to section 6.13.14.2.
1518   foreach aQual = ["RO"] in {
1519     foreach imgTy = [Image2d] in {
1520       foreach name = ["read_imagef"] in {
1521         def : Builtin<name, [VectorType<Float, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, Float], Attr.Pure>;
1522         def : Builtin<name, [VectorType<Float, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, VectorType<Float, 2>, VectorType<Float, 2>], Attr.Pure>;
1523       }
1524       foreach name = ["read_imagei"] in {
1525         def : Builtin<name, [VectorType<Int, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, Float], Attr.Pure>;
1526         def : Builtin<name, [VectorType<Int, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, VectorType<Float, 2>, VectorType<Float, 2>], Attr.Pure>;
1527       }
1528       foreach name = ["read_imageui"] in {
1529         def : Builtin<name, [VectorType<UInt, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, Float], Attr.Pure>;
1530         def : Builtin<name, [VectorType<UInt, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, VectorType<Float, 2>, VectorType<Float, 2>], Attr.Pure>;
1531       }
1532     }
1533     foreach imgTy = [Image2dDepth] in {
1534       foreach name = ["read_imagef"] in {
1535         def : Builtin<name, [Float, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, Float], Attr.Pure>;
1536         def : Builtin<name, [Float, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, VectorType<Float, 2>, VectorType<Float, 2>], Attr.Pure>;
1537       }
1538     }
1539     foreach imgTy = [Image1d] in {
1540       foreach name = ["read_imagef"] in {
1541         def : Builtin<name, [VectorType<Float, 4>, ImageType<imgTy, aQual>, Sampler, Float, Float], Attr.Pure>;
1542         def : Builtin<name, [VectorType<Float, 4>, ImageType<imgTy, aQual>, Sampler, Float, Float, Float], Attr.Pure>;
1543       }
1544       foreach name = ["read_imagei"] in {
1545         def : Builtin<name, [VectorType<Int, 4>, ImageType<imgTy, aQual>, Sampler, Float, Float], Attr.Pure>;
1546         def : Builtin<name, [VectorType<Int, 4>, ImageType<imgTy, aQual>, Sampler, Float, Float, Float], Attr.Pure>;
1547       }
1548       foreach name = ["read_imageui"] in {
1549         def : Builtin<name, [VectorType<UInt, 4>, ImageType<imgTy, aQual>, Sampler, Float, Float], Attr.Pure>;
1550         def : Builtin<name, [VectorType<UInt, 4>, ImageType<imgTy, aQual>, Sampler, Float, Float, Float], Attr.Pure>;
1551       }
1552     }
1553     foreach imgTy = [Image3d] in {
1554       foreach name = ["read_imagef"] in {
1555         def : Builtin<name, [VectorType<Float, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, VectorType<Float, 4>, VectorType<Float, 4>], Attr.Pure>;
1556         def : Builtin<name, [VectorType<Float, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, Float], Attr.Pure>;
1557       }
1558       foreach name = ["read_imagei"] in {
1559         def : Builtin<name, [VectorType<Int, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, VectorType<Float, 4>, VectorType<Float, 4>], Attr.Pure>;
1560         def : Builtin<name, [VectorType<Int, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, Float], Attr.Pure>;
1561       }
1562       foreach name = ["read_imageui"] in {
1563         def : Builtin<name, [VectorType<UInt, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, VectorType<Float, 4>, VectorType<Float, 4>], Attr.Pure>;
1564         def : Builtin<name, [VectorType<UInt, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, Float], Attr.Pure>;
1565       }
1566     }
1567     foreach imgTy = [Image1dArray] in {
1568       foreach name = ["read_imagef"] in {
1569         def : Builtin<name, [VectorType<Float, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, Float], Attr.Pure>;
1570         def : Builtin<name, [VectorType<Float, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, Float, Float], Attr.Pure>;
1571       }
1572       foreach name = ["read_imagei"] in {
1573         def : Builtin<name, [VectorType<Int, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, Float], Attr.Pure>;
1574         def : Builtin<name, [VectorType<Int, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, Float, Float], Attr.Pure>;
1575       }
1576       foreach name = ["read_imageui"] in {
1577         def : Builtin<name, [VectorType<UInt, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, Float], Attr.Pure>;
1578         def : Builtin<name, [VectorType<UInt, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, Float, Float], Attr.Pure>;
1579       }
1580     }
1581     foreach imgTy = [Image2dArray] in {
1582       foreach name = ["read_imagef"] in {
1583         def : Builtin<name, [VectorType<Float, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, Float], Attr.Pure>;
1584         def : Builtin<name, [VectorType<Float, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, VectorType<Float, 2>, VectorType<Float, 2>], Attr.Pure>;
1585       }
1586       foreach name = ["read_imagei"] in {
1587         def : Builtin<name, [VectorType<Int, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, Float], Attr.Pure>;
1588         def : Builtin<name, [VectorType<Int, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, VectorType<Float, 2>, VectorType<Float, 2>], Attr.Pure>;
1589       }
1590       foreach name = ["read_imageui"] in {
1591         def : Builtin<name, [VectorType<UInt, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, Float], Attr.Pure>;
1592         def : Builtin<name, [VectorType<UInt, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, VectorType<Float, 2>, VectorType<Float, 2>], Attr.Pure>;
1593       }
1594     }
1595     foreach imgTy = [Image2dArrayDepth] in {
1596       foreach name = ["read_imagef"] in {
1597         def : Builtin<name, [Float, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, Float], Attr.Pure>;
1598         def : Builtin<name, [Float, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, VectorType<Float, 2>, VectorType<Float, 2>], Attr.Pure>;
1599       }
1600     }
1601   }
1602 }
1603
1604 // Added to section 6.13.14.5
1605 multiclass ImageQueryNumMipLevels<string aQual> {
1606   foreach imgTy = [Image1d, Image2d, Image3d, Image1dArray, Image2dArray, Image2dDepth, Image2dArrayDepth] in {
1607     def : Builtin<"get_image_num_mip_levels", [Int, ImageType<imgTy, aQual>]>;
1608   }
1609 }
1610
1611 let Extension = FuncExtKhrMipmapImage in {
1612   defm : ImageQueryNumMipLevels<"RO">;
1613   defm : ImageQueryNumMipLevels<"WO">;
1614   defm : ImageQueryNumMipLevels<"RW">;
1615 }
1616
1617 // Write functions are enabled using a separate extension.
1618 let Extension = FuncExtKhrMipmapImageWrites in {
1619   // Added to section 6.13.14.4.
1620   foreach aQual = ["WO"] in {
1621     foreach imgTy = [Image2d] in {
1622       def : Builtin<"write_imagef", [Void, ImageType<imgTy, aQual>, VectorType<Int, 2>, Int, VectorType<Float, 4>]>;
1623       def : Builtin<"write_imagei", [Void, ImageType<imgTy, aQual>, VectorType<Int, 2>, Int, VectorType<Int, 4>]>;
1624       def : Builtin<"write_imageui", [Void, ImageType<imgTy, aQual>, VectorType<Int, 2>, Int, VectorType<UInt, 4>]>;
1625     }
1626     def : Builtin<"write_imagef", [Void, ImageType<Image2dDepth, aQual>, VectorType<Int, 2>, Int, Float]>;
1627     foreach imgTy = [Image1d] in {
1628       def : Builtin<"write_imagef", [Void, ImageType<imgTy, aQual>, Int, Int, VectorType<Float, 4>]>;
1629       def : Builtin<"write_imagei", [Void, ImageType<imgTy, aQual>, Int, Int, VectorType<Int, 4>]>;
1630       def : Builtin<"write_imageui", [Void, ImageType<imgTy, aQual>, Int, Int, VectorType<UInt, 4>]>;
1631     }
1632     foreach imgTy = [Image1dArray] in {
1633       def : Builtin<"write_imagef", [Void, ImageType<imgTy, aQual>, VectorType<Int, 2>, Int, VectorType<Float, 4>]>;
1634       def : Builtin<"write_imagei", [Void, ImageType<imgTy, aQual>, VectorType<Int, 2>, Int, VectorType<Int, 4>]>;
1635       def : Builtin<"write_imageui", [Void, ImageType<imgTy, aQual>, VectorType<Int, 2>, Int, VectorType<UInt, 4>]>;
1636     }
1637     foreach imgTy = [Image2dArray] in {
1638       def : Builtin<"write_imagef", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int, VectorType<Float, 4>]>;
1639       def : Builtin<"write_imagei", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int, VectorType<Int, 4>]>;
1640       def : Builtin<"write_imageui", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int, VectorType<UInt, 4>]>;
1641     }
1642     def : Builtin<"write_imagef", [Void, ImageType<Image2dArrayDepth, aQual>, VectorType<Int, 4>, Int, Float]>;
1643     foreach imgTy = [Image3d] in {
1644       def : Builtin<"write_imagef", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int, VectorType<Float, 4>]>;
1645       def : Builtin<"write_imagei", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int, VectorType<Int, 4>]>;
1646       def : Builtin<"write_imageui", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int, VectorType<UInt, 4>]>;
1647     }
1648   }
1649 }
1650
1651 //--------------------------------------------------------------------
1652 // OpenCL Extension v2.0 s18.3 - Creating OpenCL Memory Objects from OpenGL MSAA Textures
1653 // --- Table 6.13.14.3 ---
1654 multiclass ImageReadMsaa<string aQual> {
1655   foreach imgTy = [Image2dMsaa] in {
1656     def : Builtin<"read_imagef", [VectorType<Float, 4>, ImageType<imgTy, aQual>, VectorType<Int, 2>, Int], Attr.Pure>;
1657     def : Builtin<"read_imagei", [VectorType<Int, 4>, ImageType<imgTy, aQual>, VectorType<Int, 2>, Int], Attr.Pure>;
1658     def : Builtin<"read_imageui", [VectorType<UInt, 4>, ImageType<imgTy, aQual>, VectorType<Int, 2>, Int], Attr.Pure>;
1659   }
1660   foreach imgTy = [Image2dArrayMsaa] in {
1661     def : Builtin<"read_imagef", [VectorType<Float, 4>, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int], Attr.Pure>;
1662     def : Builtin<"read_imagei", [VectorType<Int, 4>, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int], Attr.Pure>;
1663     def : Builtin<"read_imageui", [VectorType<UInt, 4>, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int], Attr.Pure>;
1664   }
1665   foreach name = ["read_imagef"] in {
1666     def : Builtin<name, [Float, ImageType<Image2dMsaaDepth, aQual>, VectorType<Int, 2>, Int], Attr.Pure>;
1667     def : Builtin<name, [Float, ImageType<Image2dArrayMsaaDepth, aQual>, VectorType<Int, 4>, Int], Attr.Pure>;
1668   }
1669 }
1670
1671 // --- Table 6.13.14.5 ---
1672 multiclass ImageQueryMsaa<string aQual> {
1673   foreach imgTy = [Image2dMsaa, Image2dArrayMsaa, Image2dMsaaDepth, Image2dArrayMsaaDepth] in {
1674     foreach name = ["get_image_width", "get_image_height",
1675                     "get_image_channel_data_type", "get_image_channel_order",
1676                     "get_image_num_samples"] in {
1677       def : Builtin<name, [Int, ImageType<imgTy, aQual>], Attr.Const>;
1678     }
1679     def : Builtin<"get_image_dim", [VectorType<Int, 2>, ImageType<imgTy, aQual>], Attr.Const>;
1680   }
1681   foreach imgTy = [Image2dArrayMsaa, Image2dArrayMsaaDepth] in {
1682     def : Builtin<"get_image_array_size", [Size, ImageType<imgTy, aQual>], Attr.Const>;
1683   }
1684 }
1685
1686 let Extension = FuncExtKhrGlMsaaSharing in {
1687   defm : ImageReadMsaa<"RO">;
1688   defm : ImageQueryMsaa<"RO">;
1689   defm : ImageQueryMsaa<"WO">;
1690   defm : ImageReadMsaa<"RW">;
1691   defm : ImageQueryMsaa<"RW">;
1692 }
1693
1694 //--------------------------------------------------------------------
1695 // OpenCL Extension v2.0 s28 - Subgroups
1696 // --- Table 28.2.1 ---
1697 let Extension = FuncExtKhrSubgroups in {
1698   foreach name = ["get_sub_group_size", "get_max_sub_group_size",
1699                   "get_num_sub_groups", "get_sub_group_id",
1700                   "get_sub_group_local_id"] in {
1701     def : Builtin<name, [UInt]>;
1702   }
1703   let MinVersion = CL20 in {
1704     foreach name = ["get_enqueued_num_sub_groups"] in {
1705       def : Builtin<name, [UInt]>;
1706     }
1707   }
1708 }
1709
1710 // --- Table 28.2.2 ---
1711 let Extension = FuncExtKhrSubgroups in {
1712   def : Builtin<"sub_group_barrier", [Void, MemFenceFlags], Attr.Convergent>;
1713   let MinVersion = CL20 in {
1714     def : Builtin<"sub_group_barrier", [Void, MemFenceFlags, MemoryScope], Attr.Convergent>;
1715   }
1716 }
1717
1718 // --- Table 28.2.4 ---
1719 let Extension = FuncExtKhrSubgroups in {
1720   foreach name = ["sub_group_all", "sub_group_any"] in {
1721     def : Builtin<name, [Int, Int], Attr.Convergent>;
1722   }
1723   foreach name = ["sub_group_broadcast"] in {
1724     def : Builtin<name, [IntLongFloatGenType1, IntLongFloatGenType1, UInt], Attr.Convergent>;
1725   }
1726   foreach name = ["sub_group_reduce_", "sub_group_scan_exclusive_",
1727                   "sub_group_scan_inclusive_"] in {
1728     foreach op = ["add", "min", "max"] in {
1729       def : Builtin<name # op, [IntLongFloatGenType1, IntLongFloatGenType1], Attr.Convergent>;
1730     }
1731   }
1732 }
1733
1734 // OpenCL Extension v3.0 s38 - Extended Subgroup Functions
1735
1736 // Section 38.4.1 - cl_khr_subgroup_extended_types
1737 let Extension = FuncExtKhrSubgroupExtendedTypes in {
1738   // For sub_group_broadcast, add scalar char, uchar, short, and ushort support,
1739   def : Builtin<"sub_group_broadcast", [CharShortGenType1, CharShortGenType1, UInt], Attr.Convergent>;
1740   // gentype may additionally be one of the supported built-in vector data types.
1741   def : Builtin<"sub_group_broadcast", [AGenTypeNNoScalar, AGenTypeNNoScalar, UInt], Attr.Convergent>;
1742
1743   foreach name = ["sub_group_reduce_", "sub_group_scan_exclusive_",
1744                   "sub_group_scan_inclusive_"] in {
1745     foreach op = ["add", "min", "max"] in {
1746       def : Builtin<name # op, [CharShortGenType1, CharShortGenType1], Attr.Convergent>;
1747     }
1748   }
1749 }
1750
1751 // Section 38.5.1 - cl_khr_subgroup_non_uniform_vote
1752 let Extension = FuncExtKhrSubgroupNonUniformVote in {
1753   def : Builtin<"sub_group_elect", [Int]>;
1754   def : Builtin<"sub_group_non_uniform_all", [Int, Int]>;
1755   def : Builtin<"sub_group_non_uniform_any", [Int, Int]>;
1756   def : Builtin<"sub_group_non_uniform_all_equal", [Int, AGenType1]>;
1757 }
1758
1759 // Section 38.6.1 - cl_khr_subgroup_ballot
1760 let Extension = FuncExtKhrSubgroupBallot in {
1761   def : Builtin<"sub_group_non_uniform_broadcast", [AGenTypeN, AGenTypeN, UInt]>;
1762   def : Builtin<"sub_group_broadcast_first", [AGenType1, AGenType1]>;
1763   def : Builtin<"sub_group_ballot", [VectorType<UInt, 4>, Int]>;
1764   def : Builtin<"sub_group_inverse_ballot", [Int, VectorType<UInt, 4>], Attr.Const>;
1765   def : Builtin<"sub_group_ballot_bit_extract", [Int, VectorType<UInt, 4>, UInt], Attr.Const>;
1766   def : Builtin<"sub_group_ballot_bit_count", [UInt, VectorType<UInt, 4>], Attr.Const>;
1767   def : Builtin<"sub_group_ballot_inclusive_scan", [UInt, VectorType<UInt, 4>]>;
1768   def : Builtin<"sub_group_ballot_exclusive_scan", [UInt, VectorType<UInt, 4>]>;
1769   def : Builtin<"sub_group_ballot_find_lsb", [UInt, VectorType<UInt, 4>]>;
1770   def : Builtin<"sub_group_ballot_find_msb", [UInt, VectorType<UInt, 4>]>;
1771
1772   foreach op = ["eq", "ge", "gt", "le", "lt"] in {
1773     def : Builtin<"get_sub_group_" # op # "_mask", [VectorType<UInt, 4>], Attr.Const>;
1774   }
1775 }
1776
1777 // Section 38.7.1 - cl_khr_subgroup_non_uniform_arithmetic
1778 let Extension = FuncExtKhrSubgroupNonUniformArithmetic in {
1779   foreach name = ["reduce_", "scan_exclusive_", "scan_inclusive_"] in {
1780     foreach op = ["add", "min", "max", "mul"] in {
1781       def : Builtin<"sub_group_non_uniform_" # name # op, [AGenType1, AGenType1]>;
1782     }
1783     foreach op = ["and", "or", "xor"] in {
1784       def : Builtin<"sub_group_non_uniform_" # name # op, [AIGenType1, AIGenType1]>;
1785     }
1786     foreach op = ["and", "or", "xor"] in {
1787       def : Builtin<"sub_group_non_uniform_" # name # "logical_" # op, [Int, Int]>;
1788     }
1789   }
1790 }
1791
1792 // Section 38.8.1 - cl_khr_subgroup_shuffle
1793 let Extension = FuncExtKhrSubgroupShuffle in {
1794   def : Builtin<"sub_group_shuffle", [AGenType1, AGenType1, UInt]>;
1795   def : Builtin<"sub_group_shuffle_xor", [AGenType1, AGenType1, UInt]>;
1796 }
1797
1798 // Section 38.9.1 - cl_khr_subgroup_shuffle_relative
1799 let Extension = FuncExtKhrSubgroupShuffleRelative in {
1800   def : Builtin<"sub_group_shuffle_up", [AGenType1, AGenType1, UInt]>;
1801   def : Builtin<"sub_group_shuffle_down", [AGenType1, AGenType1, UInt]>;
1802 }
1803
1804 // Section 38.10.1 - cl_khr_subgroup_clustered_reduce
1805 let Extension = FuncExtKhrSubgroupClusteredReduce in {
1806   foreach op = ["add", "min", "max", "mul"] in {
1807     def : Builtin<"sub_group_clustered_reduce_" # op, [AGenType1, AGenType1, UInt]>;
1808   }
1809   foreach op = ["and", "or", "xor"] in {
1810     def : Builtin<"sub_group_clustered_reduce_" # op, [AIGenType1, AIGenType1, UInt]>;
1811   }
1812   foreach op = ["and", "or", "xor"] in {
1813     def : Builtin<"sub_group_clustered_reduce_logical_" # op, [Int, Int, UInt]>;
1814   }
1815 }
1816
1817 // Section 40.3.1 - cl_khr_extended_bit_ops
1818 let Extension = FuncExtKhrExtendedBitOps in {
1819   def : Builtin<"bitfield_insert", [AIGenTypeN, AIGenTypeN, AIGenTypeN, UInt, UInt], Attr.Const>;
1820   def : Builtin<"bitfield_extract_signed", [SGenTypeN, SGenTypeN, UInt, UInt], Attr.Const>;
1821   def : Builtin<"bitfield_extract_signed", [SGenTypeN, UGenTypeN, UInt, UInt], Attr.Const>;
1822   def : Builtin<"bitfield_extract_unsigned", [UGenTypeN, SGenTypeN, UInt, UInt], Attr.Const>;
1823   def : Builtin<"bitfield_extract_unsigned", [UGenTypeN, UGenTypeN, UInt, UInt], Attr.Const>;
1824   def : Builtin<"bit_reverse", [AIGenTypeN, AIGenTypeN], Attr.Const>;
1825 }
1826
1827 // Section 42.3 - cl_khr_integer_dot_product
1828 let Extension = FunctionExtension<"__opencl_c_integer_dot_product_input_4x8bit"> in {
1829   def : Builtin<"dot", [UInt, VectorType<UChar, 4>, VectorType<UChar, 4>], Attr.Const>;
1830   def : Builtin<"dot", [Int, VectorType<Char, 4>, VectorType<Char, 4>], Attr.Const>;
1831   def : Builtin<"dot", [Int, VectorType<UChar, 4>, VectorType<Char, 4>], Attr.Const>;
1832   def : Builtin<"dot", [Int, VectorType<Char, 4>, VectorType<UChar, 4>], Attr.Const>;
1833
1834   def : Builtin<"dot_acc_sat", [UInt, VectorType<UChar, 4>, VectorType<UChar, 4>, UInt], Attr.Const>;
1835   def : Builtin<"dot_acc_sat", [Int, VectorType<Char, 4>, VectorType<Char, 4>, Int], Attr.Const>;
1836   def : Builtin<"dot_acc_sat", [Int, VectorType<UChar, 4>, VectorType<Char, 4>, Int], Attr.Const>;
1837   def : Builtin<"dot_acc_sat", [Int, VectorType<Char, 4>, VectorType<UChar, 4>, Int], Attr.Const>;
1838 }
1839
1840 let Extension = FunctionExtension<"__opencl_c_integer_dot_product_input_4x8bit_packed"> in {
1841   def : Builtin<"dot_4x8packed_uu_uint", [UInt, UInt, UInt], Attr.Const>;
1842   def : Builtin<"dot_4x8packed_ss_int", [Int, UInt, UInt], Attr.Const>;
1843   def : Builtin<"dot_4x8packed_us_int", [Int, UInt, UInt], Attr.Const>;
1844   def : Builtin<"dot_4x8packed_su_int", [Int, UInt, UInt], Attr.Const>;
1845
1846   def : Builtin<"dot_acc_sat_4x8packed_uu_uint", [UInt, UInt, UInt, UInt], Attr.Const>;
1847   def : Builtin<"dot_acc_sat_4x8packed_ss_int", [Int, UInt, UInt, Int], Attr.Const>;
1848   def : Builtin<"dot_acc_sat_4x8packed_us_int", [Int, UInt, UInt, Int], Attr.Const>;
1849   def : Builtin<"dot_acc_sat_4x8packed_su_int", [Int, UInt, UInt, Int], Attr.Const>;
1850 }
1851
1852 // Section 48.3 - cl_khr_subgroup_rotate
1853 let Extension = FunctionExtension<"cl_khr_subgroup_rotate"> in {
1854   def : Builtin<"sub_group_rotate", [AGenType1, AGenType1, Int], Attr.Convergent>;
1855   def : Builtin<"sub_group_clustered_rotate", [AGenType1, AGenType1, Int, UInt], Attr.Convergent>;
1856 }
1857
1858 //--------------------------------------------------------------------
1859 // Arm extensions.
1860 let Extension = ArmIntegerDotProductInt8 in {
1861   foreach name = ["arm_dot"] in {
1862     def : Builtin<name, [UInt, VectorType<UChar, 4>, VectorType<UChar, 4>]>;
1863     def : Builtin<name, [Int, VectorType<Char, 4>, VectorType<Char, 4>]>;
1864   }
1865 }
1866 let Extension = ArmIntegerDotProductAccumulateInt8 in {
1867   foreach name = ["arm_dot_acc"] in {
1868     def : Builtin<name, [UInt, VectorType<UChar, 4>, VectorType<UChar, 4>, UInt]>;
1869     def : Builtin<name, [Int, VectorType<Char, 4>, VectorType<Char, 4>, Int]>;
1870   }
1871 }
1872 let Extension = ArmIntegerDotProductAccumulateInt16 in {
1873   foreach name = ["arm_dot_acc"] in {
1874     def : Builtin<name, [UInt, VectorType<UShort, 2>, VectorType<UShort, 2>, UInt]>;
1875     def : Builtin<name, [Int, VectorType<Short, 2>, VectorType<Short, 2>, Int]>;
1876   }
1877 }
1878 let Extension = ArmIntegerDotProductAccumulateSaturateInt8 in {
1879   foreach name = ["arm_dot_acc_sat"] in {
1880     def : Builtin<name, [UInt, VectorType<UChar, 4>, VectorType<UChar, 4>, UInt]>;
1881     def : Builtin<name, [Int, VectorType<Char, 4>, VectorType<Char, 4>, Int]>;
1882   }
1883 }