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                 "sqrt",
 567                 "tan", "tanh", "tanpi",
 568                 "tgamma", "trunc",
 569                 "lgamma"] in {
 570     def : Builtin<name, [FGenTypeN, FGenTypeN], Attr.Const>;
 571 }
 572 foreach name = ["nan"] in {
 573   def : Builtin<name, [GenTypeFloatVecAndScalar, GenTypeUIntVecAndScalar], Attr.Const>;
 574   def : Builtin<name, [GenTypeDoubleVecAndScalar, GenTypeULongVecAndScalar], Attr.Const>;
 575   def : Builtin<name, [GenTypeHalfVecAndScalar, GenTypeUShortVecAndScalar], Attr.Const>;
 576 }
 577
 578 // --- 2 arguments ---
 579 foreach name = ["atan2", "atan2pi", "copysign", "fdim", "fmod", "hypot",
 580                 "maxmag", "minmag", "nextafter", "pow", "powr",
 581                 "remainder"] in {
 582   def : Builtin<name, [FGenTypeN, FGenTypeN, FGenTypeN], Attr.Const>;
 583 }
 584 foreach name = ["fmax", "fmin"] in {
 585   def : Builtin<name, [FGenTypeN, FGenTypeN, FGenTypeN], Attr.Const>;
 586   def : Builtin<name, [GenTypeFloatVecNoScalar, GenTypeFloatVecNoScalar, Float], Attr.Const>;
 587   def : Builtin<name, [GenTypeDoubleVecNoScalar, GenTypeDoubleVecNoScalar, Double], Attr.Const>;
 588   def : Builtin<name, [GenTypeHalfVecNoScalar, GenTypeHalfVecNoScalar, Half], Attr.Const>;
 589 }
 590 foreach name = ["ilogb"] in {
 591   def : Builtin<name, [GenTypeIntVecAndScalar, GenTypeFloatVecAndScalar], Attr.Const>;
 592   def : Builtin<name, [GenTypeIntVecAndScalar, GenTypeDoubleVecAndScalar], Attr.Const>;
 593   def : Builtin<name, [GenTypeIntVecAndScalar, GenTypeHalfVecAndScalar], Attr.Const>;
 594 }
 595 foreach name = ["ldexp"] in {
 596   def : Builtin<name, [GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar, GenTypeIntVecAndScalar], Attr.Const>;
 597   def : Builtin<name, [GenTypeFloatVecNoScalar, GenTypeFloatVecNoScalar, Int], Attr.Const>;
 598   def : Builtin<name, [GenTypeDoubleVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeIntVecAndScalar], Attr.Const>;
 599   def : Builtin<name, [GenTypeDoubleVecNoScalar, GenTypeDoubleVecNoScalar, Int], Attr.Const>;
 600   def : Builtin<name, [GenTypeHalfVecAndScalar, GenTypeHalfVecAndScalar, GenTypeIntVecAndScalar], Attr.Const>;
 601   def : Builtin<name, [GenTypeHalfVecNoScalar, GenTypeHalfVecNoScalar, Int], Attr.Const>;
 602 }
 603 foreach name = ["pown", "rootn"] in {
 604   def : Builtin<name, [GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar, GenTypeIntVecAndScalar], Attr.Const>;
 605   def : Builtin<name, [GenTypeDoubleVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeIntVecAndScalar], Attr.Const>;
 606   def : Builtin<name, [GenTypeHalfVecAndScalar, GenTypeHalfVecAndScalar, GenTypeIntVecAndScalar], Attr.Const>;
 607 }
 608
 609 // --- 3 arguments ---
 610 foreach name = ["fma", "mad"] in {
 611   def : Builtin<name, [FGenTypeN, FGenTypeN, FGenTypeN, FGenTypeN], Attr.Const>;
 612 }
 613
 614 // The following math builtins take pointer arguments.  Which overloads are
 615 // available depends on whether the generic address space feature is enabled.
 616 multiclass MathWithPointer<list<AddressSpace> addrspaces> {
 617   foreach AS = addrspaces in {
 618     foreach name = ["fract", "modf", "sincos"] in {
 619       def : Builtin<name, [FGenTypeN, FGenTypeN, PointerType<FGenTypeN, AS>]>;
 620     }
 621     foreach name = ["frexp", "lgamma_r"] in {
 622       foreach Type = [GenTypeFloatVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeHalfVecAndScalar] in {
 623         def : Builtin<name, [Type, Type, PointerType<GenTypeIntVecAndScalar, AS>]>;
 624       }
 625     }
 626     foreach name = ["remquo"] in {
 627       foreach Type = [GenTypeFloatVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeHalfVecAndScalar] in {
 628         def : Builtin<name, [Type, Type, Type, PointerType<GenTypeIntVecAndScalar, AS>]>;
 629       }
 630     }
 631   }
 632 }
 633
 634 let Extension = FuncExtOpenCLCNamedAddressSpaceBuiltins in {
 635   defm : MathWithPointer<[GlobalAS, LocalAS, PrivateAS]>;
 636 }
 637 let Extension = FuncExtOpenCLCGenericAddressSpace in {
 638   defm : MathWithPointer<[GenericAS]>;
 639 }
 640
 641 // --- Table 9 ---
 642 foreach name = ["half_cos",
 643                 "half_exp", "half_exp2", "half_exp10",
 644                 "half_log", "half_log2", "half_log10",
 645                 "half_recip", "half_rsqrt",
 646                 "half_sin", "half_sqrt", "half_tan",
 647                 "native_cos",
 648                 "native_exp", "native_exp2", "native_exp10",
 649                 "native_log", "native_log2", "native_log10",
 650                 "native_recip", "native_rsqrt",
 651                 "native_sin", "native_sqrt", "native_tan"] in {
 652   def : Builtin<name, [GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar], Attr.Const>;
 653 }
 654 foreach name = ["half_divide", "half_powr",
 655                 "native_divide", "native_powr"] in {
 656   def : Builtin<name, [GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar], Attr.Const>;
 657 }
 658
 659 //--------------------------------------------------------------------
 660 // OpenCL v1.1 s6.11.3, v1.2 s6.12.3, v2.0 s6.13.3 - Integer Functions
 661 // --- Table 10 ---
 662 // --- 1 argument ---
 663 foreach name = ["abs"] in {
 664   def : Builtin<name, [AI2UGenTypeN, AIGenTypeN], Attr.Const>;
 665 }
 666 def : Builtin<"clz", [AIGenTypeN, AIGenTypeN], Attr.Const>;
 667 let MinVersion = CL12 in {
 668   def : Builtin<"popcount", [AIGenTypeN, AIGenTypeN], Attr.Const>;
 669 }
 670 let MinVersion = CL20 in {
 671   foreach name = ["ctz"] in {
 672     def : Builtin<name, [AIGenTypeN, AIGenTypeN], Attr.Const>;
 673   }
 674 }
 675
 676 // --- 2 arguments ---
 677 foreach name = ["abs_diff"] in {
 678   def : Builtin<name, [AI2UGenTypeN, AIGenTypeN, AIGenTypeN], Attr.Const>;
 679 }
 680 foreach name = ["add_sat", "hadd", "rhadd", "mul_hi", "rotate", "sub_sat"] in {
 681   def : Builtin<name, [AIGenTypeN, AIGenTypeN, AIGenTypeN], Attr.Const>;
 682 }
 683 foreach name = ["max", "min"] in {
 684   def : Builtin<name, [AIGenTypeN, AIGenTypeN, AIGenTypeN], Attr.Const>;
 685   def : Builtin<name, [AIGenTypeNNoScalar, AIGenTypeNNoScalar, AIGenType1], Attr.Const>;
 686 }
 687 foreach name = ["upsample"] in {
 688   def : Builtin<name, [GenTypeShortVecAndScalar, GenTypeCharVecAndScalar, GenTypeUCharVecAndScalar], Attr.Const>;
 689   def : Builtin<name, [GenTypeUShortVecAndScalar, GenTypeUCharVecAndScalar, GenTypeUCharVecAndScalar], Attr.Const>;
 690   def : Builtin<name, [GenTypeIntVecAndScalar, GenTypeShortVecAndScalar, GenTypeUShortVecAndScalar], Attr.Const>;
 691   def : Builtin<name, [GenTypeUIntVecAndScalar, GenTypeUShortVecAndScalar, GenTypeUShortVecAndScalar], Attr.Const>;
 692   def : Builtin<name, [GenTypeLongVecAndScalar, GenTypeIntVecAndScalar, GenTypeUIntVecAndScalar], Attr.Const>;
 693   def : Builtin<name, [GenTypeULongVecAndScalar, GenTypeUIntVecAndScalar, GenTypeUIntVecAndScalar], Attr.Const>;
 694 }
 695
 696 // --- 3 arguments ---
 697 foreach name = ["clamp"] in {
 698   def : Builtin<name, [AIGenTypeN, AIGenTypeN, AIGenTypeN, AIGenTypeN], Attr.Const>;
 699   def : Builtin<name, [AIGenTypeNNoScalar, AIGenTypeNNoScalar, AIGenType1, AIGenType1], Attr.Const>;
 700 }
 701 foreach name = ["mad_hi", "mad_sat"] in {
 702   def : Builtin<name, [AIGenTypeN, AIGenTypeN, AIGenTypeN, AIGenTypeN], Attr.Const>;
 703 }
 704
 705 // --- Table 11 ---
 706 foreach name = ["mad24"] in {
 707   def : Builtin<name, [GenTypeIntVecAndScalar, GenTypeIntVecAndScalar, GenTypeIntVecAndScalar, GenTypeIntVecAndScalar], Attr.Const>;
 708   def : Builtin<name, [GenTypeUIntVecAndScalar, GenTypeUIntVecAndScalar, GenTypeUIntVecAndScalar, GenTypeUIntVecAndScalar], Attr.Const>;
 709 }
 710 foreach name = ["mul24"] in {
 711   def : Builtin<name, [GenTypeIntVecAndScalar, GenTypeIntVecAndScalar, GenTypeIntVecAndScalar], Attr.Const>;
 712   def : Builtin<name, [GenTypeUIntVecAndScalar, GenTypeUIntVecAndScalar, GenTypeUIntVecAndScalar], Attr.Const>;
 713 }
 714
 715 //--------------------------------------------------------------------
 716 // OpenCL v1.1 s6.11.4, v1.2 s6.12.4, v2.0 s6.13.4 - Common Functions
 717 // OpenCL Extension v2.0 s5.1.3 and s6.1.3 - Common Functions
 718 // --- Table 12 ---
 719 // --- 1 argument ---
 720 foreach name = ["degrees", "radians", "sign"] in {
 721   def : Builtin<name, [FGenTypeN, FGenTypeN], Attr.Const>;
 722 }
 723
 724 // --- 2 arguments ---
 725 foreach name = ["max", "min"] in {
 726   def : Builtin<name, [FGenTypeN, FGenTypeN, FGenTypeN], Attr.Const>;
 727   def : Builtin<name, [GenTypeFloatVecNoScalar, GenTypeFloatVecNoScalar, Float], Attr.Const>;
 728   def : Builtin<name, [GenTypeDoubleVecNoScalar, GenTypeDoubleVecNoScalar, Double], Attr.Const>;
 729   def : Builtin<name, [GenTypeHalfVecNoScalar, GenTypeHalfVecNoScalar, Half], Attr.Const>;
 730 }
 731 foreach name = ["step"] in {
 732   def : Builtin<name, [FGenTypeN, FGenTypeN, FGenTypeN], Attr.Const>;
 733   def : Builtin<name, [GenTypeFloatVecNoScalar, Float, GenTypeFloatVecNoScalar], Attr.Const>;
 734   def : Builtin<name, [GenTypeDoubleVecNoScalar, Double, GenTypeDoubleVecNoScalar], Attr.Const>;
 735   def : Builtin<name, [GenTypeHalfVecNoScalar, Half, GenTypeHalfVecNoScalar], Attr.Const>;
 736 }
 737
 738 // --- 3 arguments ---
 739 foreach name = ["clamp"] in {
 740   def : Builtin<name, [FGenTypeN, FGenTypeN, FGenTypeN, FGenTypeN], Attr.Const>;
 741   def : Builtin<name, [GenTypeFloatVecNoScalar, GenTypeFloatVecNoScalar, Float, Float], Attr.Const>;
 742   def : Builtin<name, [GenTypeDoubleVecNoScalar, GenTypeDoubleVecNoScalar, Double, Double], Attr.Const>;
 743   def : Builtin<name, [GenTypeHalfVecNoScalar, GenTypeHalfVecNoScalar, Half, Half], Attr.Const>;
 744 }
 745 foreach name = ["mix"] in {
 746   def : Builtin<name, [FGenTypeN, FGenTypeN, FGenTypeN, FGenTypeN], Attr.Const>;
 747   def : Builtin<name, [GenTypeFloatVecNoScalar, GenTypeFloatVecNoScalar, GenTypeFloatVecNoScalar, Float], Attr.Const>;
 748   def : Builtin<name, [GenTypeDoubleVecNoScalar, GenTypeDoubleVecNoScalar, GenTypeDoubleVecNoScalar, Double], Attr.Const>;
 749   def : Builtin<name, [GenTypeHalfVecNoScalar, GenTypeHalfVecNoScalar, GenTypeHalfVecNoScalar, Half], Attr.Const>;
 750 }
 751 foreach name = ["smoothstep"] in {
 752   def : Builtin<name, [FGenTypeN, FGenTypeN, FGenTypeN, FGenTypeN], Attr.Const>;
 753   def : Builtin<name, [GenTypeFloatVecNoScalar, Float, Float, GenTypeFloatVecNoScalar], Attr.Const>;
 754   def : Builtin<name, [GenTypeDoubleVecNoScalar, Double, Double, GenTypeDoubleVecNoScalar], Attr.Const>;
 755   def : Builtin<name, [GenTypeHalfVecNoScalar, Half, Half, GenTypeHalfVecNoScalar], Attr.Const>;
 756 }
 757
 758
 759 //--------------------------------------------------------------------
 760 // OpenCL v1.1 s6.11.5, v1.2 s6.12.5, v2.0 s6.13.5 - Geometric Functions
 761 // OpenCL Extension v2.0 s5.1.4 and s6.1.4 - Geometric Functions
 762 // --- Table 13 ---
 763 // --- 1 argument ---
 764 foreach name = ["length"] in {
 765   def : Builtin<name, [Float, GenTypeFloatVec1234], Attr.Const>;
 766   def : Builtin<name, [Double, GenTypeDoubleVec1234], Attr.Const>;
 767   def : Builtin<name, [Half, GenTypeHalfVec1234], Attr.Const>;
 768 }
 769 foreach name = ["normalize"] in {
 770   def : Builtin<name, [GenTypeFloatVec1234, GenTypeFloatVec1234], Attr.Const>;
 771   def : Builtin<name, [GenTypeDoubleVec1234, GenTypeDoubleVec1234], Attr.Const>;
 772   def : Builtin<name, [GenTypeHalfVec1234, GenTypeHalfVec1234], Attr.Const>;
 773 }
 774 foreach name = ["fast_length"] in {
 775   def : Builtin<name, [Float, GenTypeFloatVec1234], Attr.Const>;
 776 }
 777 foreach name = ["fast_normalize"] in {
 778   def : Builtin<name, [GenTypeFloatVec1234, GenTypeFloatVec1234], Attr.Const>;
 779 }
 780
 781 // --- 2 arguments ---
 782 foreach name = ["cross"] in {
 783   foreach VSize = [3, 4] in {
 784     def : Builtin<name, [VectorType<Float, VSize>, VectorType<Float, VSize>, VectorType<Float, VSize>], Attr.Const>;
 785     def : Builtin<name, [VectorType<Double, VSize>, VectorType<Double, VSize>, VectorType<Double, VSize>], Attr.Const>;
 786     def : Builtin<name, [VectorType<Half, VSize>, VectorType<Half, VSize>, VectorType<Half, VSize>], Attr.Const>;
 787   }
 788 }
 789 foreach name = ["dot", "distance"] in {
 790   def : Builtin<name, [Float, GenTypeFloatVec1234, GenTypeFloatVec1234], Attr.Const>;
 791   def : Builtin<name, [Double, GenTypeDoubleVec1234, GenTypeDoubleVec1234], Attr.Const>;
 792   def : Builtin<name, [Half, GenTypeHalfVec1234, GenTypeHalfVec1234], Attr.Const>;
 793 }
 794 foreach name = ["fast_distance"] in {
 795   def : Builtin<name, [Float, GenTypeFloatVec1234, GenTypeFloatVec1234], Attr.Const>;
 796 }
 797
 798
 799 //--------------------------------------------------------------------
 800 // OpenCL v1.1 s6.11.6, v1.2 s6.12.6, v2.0 s6.13.6 - Relational Functions
 801 // OpenCL Extension v2.0 s5.1.5 and s6.1.5 - Relational Functions
 802 // --- Table 14 ---
 803 // --- 1 argument ---
 804 foreach name = ["isfinite", "isinf", "isnan", "isnormal", "signbit"] in {
 805   def : Builtin<name, [GenTypeIntVecAndScalar, GenTypeFloatVecAndScalar], Attr.Const>;
 806   def : Builtin<name, [Int, Double], Attr.Const>;
 807   def : Builtin<name, [GenTypeLongVecNoScalar, GenTypeDoubleVecNoScalar], Attr.Const>;
 808   def : Builtin<name, [Int, Half], Attr.Const>;
 809   def : Builtin<name, [GenTypeShortVecNoScalar, GenTypeHalfVecNoScalar], Attr.Const>;
 810 }
 811 foreach name = ["any", "all"] in {
 812   def : Builtin<name, [Int, SGenTypeN], Attr.Const>;
 813 }
 814
 815 // --- 2 arguments ---
 816 foreach name = ["isequal", "isnotequal", "isgreater", "isgreaterequal",
 817                 "isless", "islessequal", "islessgreater", "isordered",
 818                 "isunordered"] in {
 819   def : Builtin<name, [GenTypeIntVecAndScalar, GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar], Attr.Const>;
 820   def : Builtin<name, [Int, Double, Double], Attr.Const>;
 821   def : Builtin<name, [GenTypeLongVecNoScalar, GenTypeDoubleVecNoScalar, GenTypeDoubleVecNoScalar], Attr.Const>;
 822   def : Builtin<name, [Int, Half, Half], Attr.Const>;
 823   def : Builtin<name, [GenTypeShortVecNoScalar, GenTypeHalfVecNoScalar, GenTypeHalfVecNoScalar], Attr.Const>;
 824 }
 825
 826 // --- 3 arguments ---
 827 foreach name = ["bitselect"] in {
 828   def : Builtin<name, [AGenTypeN, AGenTypeN, AGenTypeN, AGenTypeN], Attr.Const>;
 829 }
 830 foreach name = ["select"] in {
 831   def : Builtin<name, [SGenTypeN, SGenTypeN, SGenTypeN, SGenTypeN], Attr.Const>;
 832   def : Builtin<name, [SGenTypeN, SGenTypeN, SGenTypeN, UGenTypeN], Attr.Const>;
 833   def : Builtin<name, [UGenTypeN, UGenTypeN, UGenTypeN, UGenTypeN], Attr.Const>;
 834   def : Builtin<name, [UGenTypeN, UGenTypeN, UGenTypeN, SGenTypeN], Attr.Const>;
 835   def : Builtin<name, [GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar, GenTypeIntVecAndScalar], Attr.Const>;
 836   def : Builtin<name, [GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar, GenTypeFloatVecAndScalar, GenTypeUIntVecAndScalar], Attr.Const>;
 837   def : Builtin<name, [GenTypeDoubleVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeLongVecAndScalar], Attr.Const>;
 838   def : Builtin<name, [GenTypeDoubleVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeULongVecAndScalar], Attr.Const>;
 839   def : Builtin<name, [GenTypeHalfVecAndScalar, GenTypeHalfVecAndScalar, GenTypeHalfVecAndScalar, GenTypeShortVecAndScalar], Attr.Const>;
 840   def : Builtin<name, [GenTypeHalfVecAndScalar, GenTypeHalfVecAndScalar, GenTypeHalfVecAndScalar, GenTypeUShortVecAndScalar], Attr.Const>;
 841 }
 842
 843
 844 //--------------------------------------------------------------------
 845 // OpenCL v1.1 s6.11.7, v1.2 s6.12.7, v2.0 s6.13.7 - Vector Data Load and Store Functions
 846 // 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
 847 // --- Table 15 ---
 848 multiclass VloadVstore<list<AddressSpace> addrspaces, bit defStores> {
 849   foreach AS = addrspaces in {
 850     foreach VSize = [2, 3, 4, 8, 16] in {
 851       foreach name = ["vload" # VSize] in {
 852         def : Builtin<name, [VectorType<Char, VSize>, Size, PointerType<ConstType<Char>, AS>], Attr.Pure>;
 853         def : Builtin<name, [VectorType<UChar, VSize>, Size, PointerType<ConstType<UChar>, AS>], Attr.Pure>;
 854         def : Builtin<name, [VectorType<Short, VSize>, Size, PointerType<ConstType<Short>, AS>], Attr.Pure>;
 855         def : Builtin<name, [VectorType<UShort, VSize>, Size, PointerType<ConstType<UShort>, AS>], Attr.Pure>;
 856         def : Builtin<name, [VectorType<Int, VSize>, Size, PointerType<ConstType<Int>, AS>], Attr.Pure>;
 857         def : Builtin<name, [VectorType<UInt, VSize>, Size, PointerType<ConstType<UInt>, AS>], Attr.Pure>;
 858         def : Builtin<name, [VectorType<Long, VSize>, Size, PointerType<ConstType<Long>, AS>], Attr.Pure>;
 859         def : Builtin<name, [VectorType<ULong, VSize>, Size, PointerType<ConstType<ULong>, AS>], Attr.Pure>;
 860         def : Builtin<name, [VectorType<Float, VSize>, Size, PointerType<ConstType<Float>, AS>], Attr.Pure>;
 861         def : Builtin<name, [VectorType<Double, VSize>, Size, PointerType<ConstType<Double>, AS>], Attr.Pure>;
 862         def : Builtin<name, [VectorType<Half, VSize>, Size, PointerType<ConstType<Half>, AS>], Attr.Pure>;
 863       }
 864       if defStores then {
 865         foreach name = ["vstore" # VSize] in {
 866           def : Builtin<name, [Void, VectorType<Char, VSize>, Size, PointerType<Char, AS>]>;
 867           def : Builtin<name, [Void, VectorType<UChar, VSize>, Size, PointerType<UChar, AS>]>;
 868           def : Builtin<name, [Void, VectorType<Short, VSize>, Size, PointerType<Short, AS>]>;
 869           def : Builtin<name, [Void, VectorType<UShort, VSize>, Size, PointerType<UShort, AS>]>;
 870           def : Builtin<name, [Void, VectorType<Int, VSize>, Size, PointerType<Int, AS>]>;
 871           def : Builtin<name, [Void, VectorType<UInt, VSize>, Size, PointerType<UInt, AS>]>;
 872           def : Builtin<name, [Void, VectorType<Long, VSize>, Size, PointerType<Long, AS>]>;
 873           def : Builtin<name, [Void, VectorType<ULong, VSize>, Size, PointerType<ULong, AS>]>;
 874           def : Builtin<name, [Void, VectorType<Float, VSize>, Size, PointerType<Float, AS>]>;
 875           def : Builtin<name, [Void, VectorType<Double, VSize>, Size, PointerType<Double, AS>]>;
 876           def : Builtin<name, [Void, VectorType<Half, VSize>, Size, PointerType<Half, AS>]>;
 877         }
 878       }
 879     }
 880   }
 881 }
 882
 883 let Extension = FuncExtOpenCLCNamedAddressSpaceBuiltins in {
 884   defm : VloadVstore<[GlobalAS, LocalAS, PrivateAS], 1>;
 885 }
 886 let Extension = FuncExtOpenCLCGenericAddressSpace in {
 887   defm : VloadVstore<[GenericAS], 1>;
 888 }
 889 // vload with constant address space is available regardless of version.
 890 defm : VloadVstore<[ConstantAS], 0>;
 891
 892 multiclass VloadVstoreHalf<list<AddressSpace> addrspaces, bit defStores> {
 893   foreach AS = addrspaces in {
 894     def : Builtin<"vload_half", [Float, Size, !cast<Type>("HalfPtrConst" # AS)], Attr.Pure>;
 895     foreach VSize = [2, 3, 4, 8, 16] in {
 896       foreach name = ["vload_half" # VSize, "vloada_half" # VSize] in {
 897         def : Builtin<name, [VectorType<Float, VSize>, Size, !cast<Type>("HalfPtrConst" # AS)], Attr.Pure>;
 898       }
 899     }
 900     if defStores then {
 901       foreach rnd = ["", "_rte", "_rtz", "_rtp", "_rtn"] in {
 902         foreach name = ["vstore_half" # rnd] in {
 903           def : Builtin<name, [Void, Float, Size, !cast<Type>("HalfPtr" # AS)]>;
 904           def : Builtin<name, [Void, Double, Size, !cast<Type>("HalfPtr" # AS)]>;
 905         }
 906         foreach VSize = [2, 3, 4, 8, 16] in {
 907           foreach name = ["vstore_half" # VSize # rnd, "vstorea_half" # VSize # rnd] in {
 908             def : Builtin<name, [Void, VectorType<Float, VSize>, Size, !cast<Type>("HalfPtr" # AS)]>;
 909             def : Builtin<name, [Void, VectorType<Double, VSize>, Size, !cast<Type>("HalfPtr" # AS)]>;
 910           }
 911         }
 912       }
 913     }
 914   }
 915 }
 916
 917 let Extension = FuncExtOpenCLCNamedAddressSpaceBuiltins in {
 918   defm : VloadVstoreHalf<[GlobalAS, LocalAS, PrivateAS], 1>;
 919 }
 920 let Extension = FuncExtOpenCLCGenericAddressSpace in {
 921   defm : VloadVstoreHalf<[GenericAS], 1>;
 922 }
 923 // vload_half and vloada_half with constant address space are available regardless of version.
 924 defm : VloadVstoreHalf<[ConstantAS], 0>;
 925
 926 // OpenCL v3.0 s6.15.8 - Synchronization Functions.
 927 def : Builtin<"barrier", [Void, MemFenceFlags], Attr.Convergent>;
 928 let MinVersion = CL20 in {
 929   def : Builtin<"work_group_barrier", [Void, MemFenceFlags], Attr.Convergent>;
 930   def : Builtin<"work_group_barrier", [Void, MemFenceFlags, MemoryScope], Attr.Convergent>;
 931 }
 932
 933 // OpenCL v3.0 s6.15.9 - Legacy Explicit Memory Fence Functions.
 934 def : Builtin<"mem_fence", [Void, MemFenceFlags]>;
 935 def : Builtin<"read_mem_fence", [Void, MemFenceFlags]>;
 936 def : Builtin<"write_mem_fence", [Void, MemFenceFlags]>;
 937
 938 // OpenCL v3.0 s6.15.10 - Address Space Qualifier Functions.
 939 // to_global, to_local, to_private are declared in Builtins.def.
 940
 941 let Extension = FuncExtOpenCLCGenericAddressSpace in {
 942   // The OpenCL 3.0 specification defines these with a "gentype" argument indicating any builtin
 943   // type or user-defined type, which cannot be represented currently.  Hence we slightly diverge
 944   // by providing only the following overloads with a void pointer.
 945   def : Builtin<"get_fence", [MemFenceFlags, PointerType<Void, GenericAS>]>;
 946   def : Builtin<"get_fence", [MemFenceFlags, PointerType<ConstType<Void>, GenericAS>]>;
 947 }
 948
 949 //--------------------------------------------------------------------
 950 // 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
 951 // OpenCL Extension v2.0 s5.1.7 and s6.1.7: Async Copies from Global to Local Memory, Local to Global Memory, and Prefetch
 952 // --- Table 18 ---
 953 foreach name = ["async_work_group_copy"] in {
 954   def : Builtin<name, [Event, PointerType<AGenTypeN, LocalAS>, PointerType<ConstType<AGenTypeN>, GlobalAS>, Size, Event]>;
 955   def : Builtin<name, [Event, PointerType<AGenTypeN, GlobalAS>, PointerType<ConstType<AGenTypeN>, LocalAS>, Size, Event]>;
 956 }
 957 foreach name = ["async_work_group_strided_copy"] in {
 958   def : Builtin<name, [Event, PointerType<AGenTypeN, LocalAS>, PointerType<ConstType<AGenTypeN>, GlobalAS>, Size, Size, Event]>;
 959   def : Builtin<name, [Event, PointerType<AGenTypeN, GlobalAS>, PointerType<ConstType<AGenTypeN>, LocalAS>, Size, Size, Event]>;
 960 }
 961 foreach name = ["wait_group_events"] in {
 962   def : Builtin<name, [Void, Int, PointerType<Event, GenericAS>]>;
 963 }
 964 foreach name = ["prefetch"] in {
 965   def : Builtin<name, [Void, PointerType<ConstType<AGenTypeN>, GlobalAS>, Size]>;
 966 }
 967
 968 //--------------------------------------------------------------------
 969 // OpenCL v2.0 s6.13.11 - Atomics Functions.
 970 // Functions that use memory_order and cl_mem_fence_flags enums are not
 971 // declared here as the TableGen backend does not handle enums.
 972
 973 // OpenCL v1.0 s9.5, s9.6, s9.7 - Atomic Functions for 32-bit integers
 974 // --- Table 9.1 ---
 975 let Extension = FuncExtKhrGlobalInt32BaseAtomics in {
 976   foreach Type = [Int, UInt] in {
 977     foreach name = ["atom_add", "atom_sub", "atom_xchg"] in {
 978       def : Builtin<name, [Type, PointerType<VolatileType<Type>, GlobalAS>, Type]>;
 979     }
 980     foreach name = ["atom_inc", "atom_dec"] in {
 981       def : Builtin<name, [Type, PointerType<VolatileType<Type>, GlobalAS>]>;
 982     }
 983     foreach name = ["atom_cmpxchg"] in {
 984       def : Builtin<name, [Type, PointerType<VolatileType<Type>, GlobalAS>, Type, Type]>;
 985     }
 986   }
 987 }
 988 // --- Table 9.3 ---
 989 let Extension = FuncExtKhrLocalInt32BaseAtomics in {
 990   foreach Type = [Int, UInt] in {
 991     foreach name = ["atom_add", "atom_sub", "atom_xchg"] in {
 992       def : Builtin<name, [Type, PointerType<VolatileType<Type>, LocalAS>, Type]>;
 993     }
 994     foreach name = ["atom_inc", "atom_dec"] in {
 995       def : Builtin<name, [Type, PointerType<VolatileType<Type>, LocalAS>]>;
 996     }
 997     foreach name = ["atom_cmpxchg"] in {
 998       def : Builtin<name, [Type, PointerType<VolatileType<Type>, LocalAS>, Type, Type]>;
 999     }
1000   }
1001 }
1002 // --- Table 9.5 ---
1003 let Extension = FuncExtKhrInt64BaseAtomics in {
1004   foreach AS = [GlobalAS, LocalAS] in {
1005     foreach Type = [Long, ULong] in {
1006       foreach name = ["atom_add", "atom_sub", "atom_xchg"] in {
1007         def : Builtin<name, [Type, PointerType<VolatileType<Type>, AS>, Type]>;
1008       }
1009       foreach name = ["atom_inc", "atom_dec"] in {
1010         def : Builtin<name, [Type, PointerType<VolatileType<Type>, AS>]>;
1011       }
1012       foreach name = ["atom_cmpxchg"] in {
1013         def : Builtin<name, [Type, PointerType<VolatileType<Type>, AS>, Type, Type]>;
1014       }
1015     }
1016   }
1017 }
1018 // --- Table 9.2 ---
1019 let Extension = FuncExtKhrGlobalInt32ExtendedAtomics in {
1020   foreach Type = [Int, UInt] in {
1021     foreach name = ["atom_min", "atom_max", "atom_and",
1022                     "atom_or", "atom_xor"] in {
1023       def : Builtin<name, [Type, PointerType<VolatileType<Type>, GlobalAS>, Type]>;
1024     }
1025   }
1026 }
1027 // --- Table 9.4 ---
1028 let Extension = FuncExtKhrLocalInt32ExtendedAtomics in {
1029   foreach Type = [Int, UInt] in {
1030     foreach name = ["atom_min", "atom_max", "atom_and",
1031                     "atom_or", "atom_xor"] in {
1032       def : Builtin<name, [Type, PointerType<VolatileType<Type>, LocalAS>, Type]>;
1033     }
1034   }
1035 }
1036 // --- Table 9.6 ---
1037 let Extension = FuncExtKhrInt64ExtendedAtomics in {
1038   foreach AS = [GlobalAS, LocalAS] in {
1039     foreach Type = [Long, ULong] in {
1040       foreach name = ["atom_min", "atom_max", "atom_and",
1041                       "atom_or", "atom_xor"] in {
1042         def : Builtin<name, [Type, PointerType<VolatileType<Type>, AS>, Type]>;
1043       }
1044     }
1045   }
1046 }
1047 // OpenCL v1.1 s6.11.1, v1.2 s6.12.11 - Atomic Functions
1048 foreach AS = [GlobalAS, LocalAS] in {
1049   def : Builtin<"atomic_xchg", [Float, PointerType<VolatileType<Float>, AS>, Float]>;
1050   foreach Type = [Int, UInt] in {
1051     foreach name = ["atomic_add", "atomic_sub", "atomic_xchg",
1052                     "atomic_min", "atomic_max", "atomic_and",
1053                     "atomic_or", "atomic_xor"] in {
1054       def : Builtin<name, [Type, PointerType<VolatileType<Type>, AS>, Type]>;
1055     }
1056     foreach name = ["atomic_inc", "atomic_dec"] in {
1057       def : Builtin<name, [Type, PointerType<VolatileType<Type>, AS>]>;
1058     }
1059     foreach name = ["atomic_cmpxchg"] in {
1060       def : Builtin<name, [Type, PointerType<VolatileType<Type>, AS>, Type, Type]>;
1061     }
1062   }
1063 }
1064
1065 let Extension = FuncExtOpenCLCxx in {
1066   foreach Type = [Int, UInt] in {
1067     foreach name = ["atomic_add", "atomic_sub", "atomic_xchg",
1068                     "atomic_min", "atomic_max", "atomic_and",
1069                     "atomic_or", "atomic_xor"] in {
1070       def : Builtin<name, [Type, PointerType<VolatileType<Type>, GenericAS>, Type]>;
1071     }
1072     foreach name = ["atomic_inc", "atomic_dec"] in {
1073       def : Builtin<name, [Type, PointerType<VolatileType<Type>, GenericAS>]>;
1074     }
1075     foreach name = ["atomic_cmpxchg"] in {
1076       def : Builtin<name, [Type, PointerType<VolatileType<Type>, GenericAS>, Type, Type]>;
1077     }
1078   }
1079 }
1080
1081 // OpenCL v2.0 s6.13.11 - Atomic Functions.
1082
1083 // An atomic builtin with 2 additional _explicit variants.
1084 multiclass BuiltinAtomicExplicit<string Name, list<Type> Types, FunctionExtension BaseExt> {
1085   // Without explicit MemoryOrder or MemoryScope.
1086   let Extension = concatExtension<BaseExt, "__opencl_c_atomic_order_seq_cst __opencl_c_atomic_scope_device">.ret in {
1087     def : Builtin<Name, Types>;
1088   }
1089
1090   // With an explicit MemoryOrder argument.
1091   let Extension = concatExtension<BaseExt, "__opencl_c_atomic_scope_device">.ret in {
1092     def : Builtin<Name # "_explicit", !listconcat(Types, [MemoryOrder])>;
1093   }
1094
1095   // With explicit MemoryOrder and MemoryScope arguments.
1096   let Extension = BaseExt in {
1097     def : Builtin<Name # "_explicit", !listconcat(Types, [MemoryOrder, MemoryScope])>;
1098   }
1099 }
1100
1101 // OpenCL 2.0 atomic functions that have a pointer argument in a given address space.
1102 multiclass OpenCL2Atomics<AddressSpace addrspace, FunctionExtension BaseExt> {
1103   foreach TypePair = [[AtomicInt, Int], [AtomicUInt, UInt],
1104                       [AtomicLong, Long], [AtomicULong, ULong],
1105                       [AtomicFloat, Float], [AtomicDouble, Double]] in {
1106     let Extension = BaseExt in {
1107       def : Builtin<"atomic_init",
1108           [Void, PointerType<VolatileType<TypePair[0]>, addrspace>, TypePair[1]]>;
1109     }
1110     defm : BuiltinAtomicExplicit<"atomic_store",
1111         [Void, PointerType<VolatileType<TypePair[0]>, addrspace>, TypePair[1]], BaseExt>;
1112     defm : BuiltinAtomicExplicit<"atomic_load",
1113         [TypePair[1], PointerType<VolatileType<TypePair[0]>, addrspace>], BaseExt>;
1114     defm : BuiltinAtomicExplicit<"atomic_exchange",
1115         [TypePair[1], PointerType<VolatileType<TypePair[0]>, addrspace>, TypePair[1]], BaseExt>;
1116     foreach Variant = ["weak", "strong"] in {
1117       foreach exp_ptr_addrspace = !cond(
1118             !eq(BaseExt, FuncExtOpenCLCGenericAddressSpace): [GenericAS],
1119             !eq(BaseExt, FuncExtOpenCLCNamedAddressSpaceBuiltins): [GlobalAS, LocalAS, PrivateAS])
1120           in {
1121         let Extension = concatExtension<BaseExt, "__opencl_c_atomic_order_seq_cst __opencl_c_atomic_scope_device">.ret in {
1122           def : Builtin<"atomic_compare_exchange_" # Variant,
1123               [Bool, PointerType<VolatileType<TypePair[0]>, addrspace>,
1124                PointerType<TypePair[1], exp_ptr_addrspace>, TypePair[1]]>;
1125         }
1126         let Extension = concatExtension<BaseExt, "__opencl_c_atomic_scope_device">.ret in {
1127           def : Builtin<"atomic_compare_exchange_" # Variant # "_explicit",
1128               [Bool, PointerType<VolatileType<TypePair[0]>, addrspace>,
1129                PointerType<TypePair[1], exp_ptr_addrspace>, TypePair[1], MemoryOrder, MemoryOrder]>;
1130         }
1131         let Extension = BaseExt in {
1132           def : Builtin<"atomic_compare_exchange_" # Variant # "_explicit",
1133               [Bool, PointerType<VolatileType<TypePair[0]>, addrspace>,
1134                PointerType<TypePair[1], exp_ptr_addrspace>, TypePair[1], MemoryOrder, MemoryOrder, MemoryScope]>;
1135         }
1136       }
1137     }
1138   }
1139
1140   foreach TypePair = [[AtomicInt, Int, Int], [AtomicUInt, UInt, UInt],
1141                       [AtomicLong, Long, Long], [AtomicULong, ULong, ULong],
1142                       [AtomicUIntPtr, UIntPtr, PtrDiff]] in {
1143     foreach ModOp = ["add", "sub"] in {
1144       defm : BuiltinAtomicExplicit<"atomic_fetch_" # ModOp,
1145           [TypePair[1], PointerType<VolatileType<TypePair[0]>, addrspace>, TypePair[2]], BaseExt>;
1146     }
1147   }
1148   foreach TypePair = [[AtomicInt, Int, Int], [AtomicUInt, UInt, UInt],
1149                       [AtomicLong, Long, Long], [AtomicULong, ULong, ULong]] in {
1150     foreach ModOp = ["or", "xor", "and", "min", "max"] in {
1151       defm : BuiltinAtomicExplicit<"atomic_fetch_" # ModOp,
1152           [TypePair[1], PointerType<VolatileType<TypePair[0]>, addrspace>, TypePair[2]], BaseExt>;
1153     }
1154   }
1155
1156   defm : BuiltinAtomicExplicit<"atomic_flag_clear",
1157       [Void, PointerType<VolatileType<AtomicFlag>, addrspace>], BaseExt>;
1158
1159   defm : BuiltinAtomicExplicit<"atomic_flag_test_and_set",
1160       [Bool, PointerType<VolatileType<AtomicFlag>, addrspace>], BaseExt>;
1161 }
1162
1163 let MinVersion = CL20 in {
1164   def : Builtin<"atomic_work_item_fence", [Void, MemFenceFlags, MemoryOrder, MemoryScope]>;
1165
1166   defm : OpenCL2Atomics<GenericAS, FuncExtOpenCLCGenericAddressSpace>;
1167   defm : OpenCL2Atomics<GlobalAS, FuncExtOpenCLCNamedAddressSpaceBuiltins>;
1168   defm : OpenCL2Atomics<LocalAS, FuncExtOpenCLCNamedAddressSpaceBuiltins>;
1169 }
1170
1171 // The functionality added by cl_ext_float_atomics extension
1172 let MinVersion = CL20 in {
1173   foreach addrspace = [GlobalAS, LocalAS, GenericAS] in {
1174     defvar extension_fp16 = !cast<FunctionExtension>("FuncExtFloatAtomicsFp16" # addrspace # "LoadStore");
1175
1176     defm : BuiltinAtomicExplicit<"atomic_store",
1177         [Void, PointerType<VolatileType<AtomicHalf>, addrspace>, AtomicHalf], extension_fp16>;
1178     defm : BuiltinAtomicExplicit<"atomic_load",
1179         [Half, PointerType<VolatileType<AtomicHalf>, addrspace>], extension_fp16>;
1180     defm : BuiltinAtomicExplicit<"atomic_exchange",
1181         [Half, PointerType<VolatileType<AtomicHalf>, addrspace>, Half], extension_fp16>;
1182
1183     foreach ModOp = ["add", "sub"] in {
1184       defvar extension_fp16 = !cast<FunctionExtension>("FuncExtFloatAtomicsFp16" # addrspace # "Add");
1185       defvar extension_fp32 = !cast<FunctionExtension>("FuncExtFloatAtomicsFp32" # addrspace # "Add");
1186       defvar extension_fp64 = !cast<FunctionExtension>("FuncExtFloatAtomicsFp64" # addrspace # "Add");
1187
1188       defm : BuiltinAtomicExplicit<"atomic_fetch_" # ModOp,
1189           [Half, PointerType<VolatileType<AtomicHalf>, addrspace>, Half], extension_fp16>;
1190       defm : BuiltinAtomicExplicit<"atomic_fetch_" # ModOp,
1191           [Float, PointerType<VolatileType<AtomicFloat>, addrspace>, Float], extension_fp32>;
1192       defm : BuiltinAtomicExplicit<"atomic_fetch_" # ModOp,
1193           [Double, PointerType<VolatileType<AtomicDouble>, addrspace>, Double], extension_fp64>;
1194     }
1195
1196     foreach ModOp = ["min", "max"] in {
1197       defvar extension_fp16 = !cast<FunctionExtension>("FuncExtFloatAtomicsFp16" # addrspace # "MinMax");
1198       defvar extension_fp32 = !cast<FunctionExtension>("FuncExtFloatAtomicsFp32" # addrspace # "MinMax");
1199       defvar extension_fp64 = !cast<FunctionExtension>("FuncExtFloatAtomicsFp64" # addrspace # "MinMax");
1200
1201       defm : BuiltinAtomicExplicit<"atomic_fetch_" # ModOp,
1202           [Half, PointerType<VolatileType<AtomicHalf>, addrspace>, Half], extension_fp16>;
1203       defm : BuiltinAtomicExplicit<"atomic_fetch_" # ModOp,
1204           [Float, PointerType<VolatileType<AtomicFloat>, addrspace>, Float], extension_fp32>;
1205       defm : BuiltinAtomicExplicit<"atomic_fetch_" # ModOp,
1206           [Double, PointerType<VolatileType<AtomicDouble>, addrspace>, Double], extension_fp64>;
1207     }
1208   }
1209 }
1210
1211 //--------------------------------------------------------------------
1212 // OpenCL v1.1 s6.11.12, v1.2 s6.12.12, v2.0 s6.13.12 - Miscellaneous Vector Functions
1213 // --- Table 19 ---
1214 foreach VSize1 = [2, 4, 8, 16] in {
1215   foreach VSize2 = [2, 4, 8, 16] in {
1216     foreach VecAndMaskType = [[Char, UChar], [UChar, UChar],
1217                               [Short, UShort], [UShort, UShort],
1218                               [Int, UInt], [UInt, UInt],
1219                               [Long, ULong], [ULong, ULong],
1220                               [Float, UInt], [Double, ULong], [Half, UShort]] in {
1221       def : Builtin<"shuffle", [VectorType<VecAndMaskType[0], VSize1>,
1222                                 VectorType<VecAndMaskType[0], VSize2>,
1223                                 VectorType<VecAndMaskType[1], VSize1>],
1224                                Attr.Const>;
1225     }
1226   }
1227 }
1228 foreach VSize1 = [2, 4, 8, 16] in {
1229   foreach VSize2 = [2, 4, 8, 16] in {
1230     foreach VecAndMaskType = [[Char, UChar], [UChar, UChar],
1231                               [Short, UShort], [UShort, UShort],
1232                               [Int, UInt], [UInt, UInt],
1233                               [Long, ULong], [ULong, ULong],
1234                               [Float, UInt], [Double, ULong], [Half, UShort]] in {
1235       def : Builtin<"shuffle2", [VectorType<VecAndMaskType[0], VSize1>,
1236                                  VectorType<VecAndMaskType[0], VSize2>,
1237                                  VectorType<VecAndMaskType[0], VSize2>,
1238                                  VectorType<VecAndMaskType[1], VSize1>],
1239                                 Attr.Const>;
1240     }
1241   }
1242 }
1243
1244 //--------------------------------------------------------------------
1245 // OpenCL v1.1 s6.11.3, v1.2 s6.12.14, v2.0 s6.13.14: Image Read and Write Functions
1246 // OpenCL Extension v2.0 s5.1.8 and s6.1.8: Image Read and Write Functions
1247 // --- Table 22: Image Read Functions with Samplers ---
1248 foreach imgTy = [Image1d] in {
1249   foreach coordTy = [Int, Float] in {
1250     def : Builtin<"read_imagef", [VectorType<Float, 4>, ImageType<imgTy, "RO">, Sampler, coordTy], Attr.Pure>;
1251     def : Builtin<"read_imagei", [VectorType<Int, 4>, ImageType<imgTy, "RO">, Sampler, coordTy], Attr.Pure>;
1252     def : Builtin<"read_imageui", [VectorType<UInt, 4>, ImageType<imgTy, "RO">, Sampler, coordTy], Attr.Pure>;
1253   }
1254 }
1255 foreach imgTy = [Image2d, Image1dArray] in {
1256   foreach coordTy = [Int, Float] in {
1257     def : Builtin<"read_imagef", [VectorType<Float, 4>, ImageType<imgTy, "RO">, Sampler, VectorType<coordTy, 2>], Attr.Pure>;
1258     def : Builtin<"read_imagei", [VectorType<Int, 4>, ImageType<imgTy, "RO">, Sampler, VectorType<coordTy, 2>], Attr.Pure>;
1259     def : Builtin<"read_imageui", [VectorType<UInt, 4>, ImageType<imgTy, "RO">, Sampler, VectorType<coordTy, 2>], Attr.Pure>;
1260   }
1261 }
1262 foreach imgTy = [Image3d, Image2dArray] in {
1263   foreach coordTy = [Int, Float] in {
1264     def : Builtin<"read_imagef", [VectorType<Float, 4>, ImageType<imgTy, "RO">, Sampler, VectorType<coordTy, 4>], Attr.Pure>;
1265     def : Builtin<"read_imagei", [VectorType<Int, 4>, ImageType<imgTy, "RO">, Sampler, VectorType<coordTy, 4>], Attr.Pure>;
1266     def : Builtin<"read_imageui", [VectorType<UInt, 4>, ImageType<imgTy, "RO">, Sampler, VectorType<coordTy, 4>], Attr.Pure>;
1267   }
1268 }
1269 foreach coordTy = [Int, Float] in {
1270   def : Builtin<"read_imagef", [Float, ImageType<Image2dDepth, "RO">, Sampler, VectorType<coordTy, 2>], Attr.Pure>;
1271   def : Builtin<"read_imagef", [Float, ImageType<Image2dArrayDepth, "RO">, Sampler, VectorType<coordTy, 4>], Attr.Pure>;
1272 }
1273
1274 // --- Table 23: Sampler-less Read Functions ---
1275 multiclass ImageReadSamplerless<string aQual> {
1276   foreach imgTy = [Image2d, Image1dArray] in {
1277     def : Builtin<"read_imagef", [VectorType<Float, 4>, ImageType<imgTy, aQual>, VectorType<Int, 2>], Attr.Pure>;
1278     def : Builtin<"read_imagei", [VectorType<Int, 4>, ImageType<imgTy, aQual>, VectorType<Int, 2>], Attr.Pure>;
1279     def : Builtin<"read_imageui", [VectorType<UInt, 4>, ImageType<imgTy, aQual>, VectorType<Int, 2>], Attr.Pure>;
1280   }
1281   foreach imgTy = [Image3d, Image2dArray] in {
1282     def : Builtin<"read_imagef", [VectorType<Float, 4>, ImageType<imgTy, aQual>, VectorType<Int, 4>], Attr.Pure>;
1283     def : Builtin<"read_imagei", [VectorType<Int, 4>, ImageType<imgTy, aQual>, VectorType<Int, 4>], Attr.Pure>;
1284     def : Builtin<"read_imageui", [VectorType<UInt, 4>, ImageType<imgTy, aQual>, VectorType<Int, 4>], Attr.Pure>;
1285   }
1286   foreach imgTy = [Image1d, Image1dBuffer] in {
1287     def : Builtin<"read_imagef", [VectorType<Float, 4>, ImageType<imgTy, aQual>, Int], Attr.Pure>;
1288     def : Builtin<"read_imagei", [VectorType<Int, 4>, ImageType<imgTy, aQual>, Int], Attr.Pure>;
1289     def : Builtin<"read_imageui", [VectorType<UInt, 4>, ImageType<imgTy, aQual>, Int], Attr.Pure>;
1290   }
1291   def : Builtin<"read_imagef", [Float, ImageType<Image2dDepth, aQual>, VectorType<Int, 2>], Attr.Pure>;
1292   def : Builtin<"read_imagef", [Float, ImageType<Image2dArrayDepth, aQual>, VectorType<Int, 4>], Attr.Pure>;
1293 }
1294
1295 let MinVersion = CL12 in {
1296   defm : ImageReadSamplerless<"RO">;
1297   let Extension = FuncExtOpenCLCReadWriteImages in {
1298     defm : ImageReadSamplerless<"RW">;
1299   }
1300 }
1301
1302 // --- Table 24: Image Write Functions ---
1303 multiclass ImageWrite<string aQual> {
1304   foreach imgTy = [Image2d] in {
1305     def : Builtin<"write_imagef", [Void, ImageType<imgTy, aQual>, VectorType<Int, 2>, VectorType<Float, 4>]>;
1306     def : Builtin<"write_imagei", [Void, ImageType<imgTy, aQual>, VectorType<Int, 2>, VectorType<Int, 4>]>;
1307     def : Builtin<"write_imageui", [Void, ImageType<imgTy, aQual>, VectorType<Int, 2>, VectorType<UInt, 4>]>;
1308   }
1309   foreach imgTy = [Image2dArray] in {
1310     def : Builtin<"write_imagef", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, VectorType<Float, 4>]>;
1311     def : Builtin<"write_imagei", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, VectorType<Int, 4>]>;
1312     def : Builtin<"write_imageui", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, VectorType<UInt, 4>]>;
1313   }
1314   foreach imgTy = [Image1d, Image1dBuffer] in {
1315     def : Builtin<"write_imagef", [Void, ImageType<imgTy, aQual>, Int, VectorType<Float, 4>]>;
1316     def : Builtin<"write_imagei", [Void, ImageType<imgTy, aQual>, Int, VectorType<Int, 4>]>;
1317     def : Builtin<"write_imageui", [Void, ImageType<imgTy, aQual>, Int, VectorType<UInt, 4>]>;
1318   }
1319   foreach imgTy = [Image1dArray] in {
1320     def : Builtin<"write_imagef", [Void, ImageType<imgTy, aQual>, VectorType<Int, 2>, VectorType<Float, 4>]>;
1321     def : Builtin<"write_imagei", [Void, ImageType<imgTy, aQual>, VectorType<Int, 2>, VectorType<Int, 4>]>;
1322     def : Builtin<"write_imageui", [Void, ImageType<imgTy, aQual>, VectorType<Int, 2>, VectorType<UInt, 4>]>;
1323   }
1324   foreach imgTy = [Image3d] in {
1325     def : Builtin<"write_imagef", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, VectorType<Float, 4>]>;
1326     def : Builtin<"write_imagei", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, VectorType<Int, 4>]>;
1327     def : Builtin<"write_imageui", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, VectorType<UInt, 4>]>;
1328   }
1329   def : Builtin<"write_imagef", [Void, ImageType<Image2dDepth, aQual>, VectorType<Int, 2>, Float]>;
1330   def : Builtin<"write_imagef", [Void, ImageType<Image2dArrayDepth, aQual>, VectorType<Int, 4>, Float]>;
1331 }
1332
1333 defm : ImageWrite<"WO">;
1334 let Extension = FuncExtOpenCLCReadWriteImages in {
1335   defm : ImageWrite<"RW">;
1336 }
1337
1338 // --- Table 25: Image Query Functions ---
1339 multiclass ImageQuery<string aQual> {
1340   foreach imgTy = [Image1d, Image1dBuffer, Image2d, Image3d,
1341                    Image1dArray, Image2dArray, Image2dDepth,
1342                    Image2dArrayDepth] in {
1343     foreach name = ["get_image_width", "get_image_channel_data_type",
1344                     "get_image_channel_order"] in {
1345       def : Builtin<name, [Int, ImageType<imgTy, aQual>], Attr.Const>;
1346     }
1347   }
1348   foreach imgTy = [Image2d, Image3d, Image2dArray, Image2dDepth,
1349                    Image2dArrayDepth] in {
1350     def : Builtin<"get_image_height", [Int, ImageType<imgTy, aQual>], Attr.Const>;
1351   }
1352   def : Builtin<"get_image_depth", [Int, ImageType<Image3d, aQual>], Attr.Const>;
1353   foreach imgTy = [Image2d, Image2dArray, Image2dDepth,
1354                    Image2dArrayDepth] in {
1355     def : Builtin<"get_image_dim", [VectorType<Int, 2>, ImageType<imgTy, aQual>], Attr.Const>;
1356   }
1357   def : Builtin<"get_image_dim", [VectorType<Int, 4>, ImageType<Image3d, aQual>], Attr.Const>;
1358   foreach imgTy = [Image1dArray, Image2dArray, Image2dArrayDepth] in {
1359     def : Builtin<"get_image_array_size", [Size, ImageType<imgTy, aQual>], Attr.Const>;
1360   }
1361 }
1362
1363 defm : ImageQuery<"RO">;
1364 defm : ImageQuery<"WO">;
1365 let Extension = FuncExtOpenCLCReadWriteImages in {
1366   defm : ImageQuery<"RW">;
1367 }
1368
1369 // OpenCL extension v2.0 s5.1.9: Built-in Image Read Functions
1370 // --- Table 8 ---
1371 foreach aQual = ["RO"] in {
1372   foreach name = ["read_imageh"] in {
1373     foreach coordTy = [Int, Float] in {
1374       foreach imgTy = [Image2d, Image1dArray] in {
1375         def : Builtin<name, [VectorType<Half, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<coordTy, 2>], Attr.Pure>;
1376       }
1377       foreach imgTy = [Image3d, Image2dArray] in {
1378         def : Builtin<name, [VectorType<Half, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<coordTy, 4>], Attr.Pure>;
1379       }
1380       foreach imgTy = [Image1d] in {
1381         def : Builtin<name, [VectorType<Half, 4>, ImageType<imgTy, aQual>, Sampler, coordTy], Attr.Pure>;
1382       }
1383     }
1384   }
1385 }
1386 // OpenCL extension v2.0 s5.1.10: Built-in Image Sampler-less Read Functions
1387 // --- Table 9 ---
1388 let MinVersion = CL12 in {
1389   multiclass ImageReadHalf<string aQual> {
1390     foreach name = ["read_imageh"] in {
1391       foreach imgTy = [Image2d, Image1dArray] in {
1392         def : Builtin<name, [VectorType<Half, 4>, ImageType<imgTy, aQual>, VectorType<Int, 2>], Attr.Pure>;
1393       }
1394       foreach imgTy = [Image3d, Image2dArray] in {
1395         def : Builtin<name, [VectorType<Half, 4>, ImageType<imgTy, aQual>, VectorType<Int, 4>], Attr.Pure>;
1396       }
1397       foreach imgTy = [Image1d, Image1dBuffer] in {
1398         def : Builtin<name, [VectorType<Half, 4>, ImageType<imgTy, aQual>, Int], Attr.Pure>;
1399       }
1400     }
1401   }
1402   defm : ImageReadHalf<"RO">;
1403   let Extension = FuncExtOpenCLCReadWriteImages in {
1404     defm : ImageReadHalf<"RW">;
1405   }
1406 }
1407 // OpenCL extension v2.0 s5.1.11: Built-in Image Write Functions
1408 // --- Table 10 ---
1409 multiclass ImageWriteHalf<string aQual> {
1410   foreach name = ["write_imageh"] in {
1411     def : Builtin<name, [Void, ImageType<Image2d, aQual>, VectorType<Int, 2>, VectorType<Half, 4>]>;
1412     def : Builtin<name, [Void, ImageType<Image2dArray, aQual>, VectorType<Int, 4>, VectorType<Half, 4>]>;
1413     def : Builtin<name, [Void, ImageType<Image1d, aQual>, Int, VectorType<Half, 4>]>;
1414     def : Builtin<name, [Void, ImageType<Image1dBuffer, aQual>, Int, VectorType<Half, 4>]>;
1415     def : Builtin<name, [Void, ImageType<Image1dArray, aQual>, VectorType<Int, 2>, VectorType<Half, 4>]>;
1416     def : Builtin<name, [Void, ImageType<Image3d, aQual>, VectorType<Int, 4>, VectorType<Half, 4>]>;
1417   }
1418 }
1419
1420 defm : ImageWriteHalf<"WO">;
1421 let Extension = FuncExtOpenCLCReadWriteImages in {
1422   defm : ImageWriteHalf<"RW">;
1423 }
1424
1425
1426
1427 //--------------------------------------------------------------------
1428 // OpenCL v2.0 s6.13.15 - Work-group Functions
1429 // --- Table 26 ---
1430 let Extension = FuncExtOpenCLCWGCollectiveFunctions in {
1431   foreach name = ["work_group_all", "work_group_any"] in {
1432     def : Builtin<name, [Int, Int], Attr.Convergent>;
1433   }
1434   foreach name = ["work_group_broadcast"] in {
1435     def : Builtin<name, [IntLongFloatGenType1, IntLongFloatGenType1, Size], Attr.Convergent>;
1436     def : Builtin<name, [IntLongFloatGenType1, IntLongFloatGenType1, Size, Size], Attr.Convergent>;
1437     def : Builtin<name, [IntLongFloatGenType1, IntLongFloatGenType1, Size, Size, Size], Attr.Convergent>;
1438   }
1439   foreach op = ["add", "min", "max"] in {
1440     foreach name = ["work_group_reduce_", "work_group_scan_exclusive_",
1441                     "work_group_scan_inclusive_"] in {
1442       def : Builtin<name # op, [IntLongFloatGenType1, IntLongFloatGenType1], Attr.Convergent>;
1443     }
1444   }
1445 }
1446
1447
1448 //--------------------------------------------------------------------
1449 // OpenCL2.0 : 6.13.16 : Pipe Functions
1450 // --- Table 27 ---
1451 // Defined in Builtins.def
1452
1453 // --- Table 28 ---
1454 // Builtins taking pipe arguments are defined in Builtins.def
1455 let Extension = FuncExtOpenCLCPipes in {
1456   def : Builtin<"is_valid_reserve_id", [Bool, ReserveId]>;
1457 }
1458
1459 // --- Table 29 ---
1460 // Defined in Builtins.def
1461
1462
1463 //--------------------------------------------------------------------
1464 // OpenCL2.0 : 6.13.17 : Enqueuing Kernels
1465 // --- Table 30 ---
1466 // Defined in Builtins.def
1467
1468 // --- Table 32 ---
1469 // Defined in Builtins.def
1470
1471 // --- Table 33 ---
1472 let Extension = FuncExtOpenCLCDeviceEnqueue in {
1473   def : Builtin<"enqueue_marker",
1474       [Int, Queue, UInt, PointerType<ConstType<ClkEvent>, GenericAS>, PointerType<ClkEvent, GenericAS>]>;
1475
1476   // --- Table 34 ---
1477   def : Builtin<"retain_event", [Void, ClkEvent]>;
1478   def : Builtin<"release_event", [Void, ClkEvent]>;
1479   def : Builtin<"create_user_event", [ClkEvent]>;
1480   def : Builtin<"is_valid_event", [Bool, ClkEvent]>;
1481   def : Builtin<"set_user_event_status", [Void, ClkEvent, Int]>;
1482   def : Builtin<"capture_event_profiling_info",
1483       [Void, ClkEvent, ClkProfilingInfo, PointerType<Void, GlobalAS>]>;
1484
1485   // --- Table 35 ---
1486   def : Builtin<"get_default_queue", [Queue]>;
1487
1488   def : Builtin<"ndrange_1D", [NDRange, Size]>;
1489   def : Builtin<"ndrange_1D", [NDRange, Size, Size]>;
1490   def : Builtin<"ndrange_1D", [NDRange, Size, Size, Size]>;
1491   def : Builtin<"ndrange_2D", [NDRange, PointerType<ConstType<Size>, PrivateAS>]>;
1492   def : Builtin<"ndrange_2D", [NDRange, PointerType<ConstType<Size>, PrivateAS>,
1493                                         PointerType<ConstType<Size>, PrivateAS>]>;
1494   def : Builtin<"ndrange_2D", [NDRange, PointerType<ConstType<Size>, PrivateAS>,
1495                                         PointerType<ConstType<Size>, PrivateAS>,
1496                                         PointerType<ConstType<Size>, PrivateAS>]>;
1497   def : Builtin<"ndrange_3D", [NDRange, PointerType<ConstType<Size>, PrivateAS>]>;
1498   def : Builtin<"ndrange_3D", [NDRange, PointerType<ConstType<Size>, PrivateAS>,
1499                                         PointerType<ConstType<Size>, PrivateAS>]>;
1500   def : Builtin<"ndrange_3D", [NDRange, PointerType<ConstType<Size>, PrivateAS>,
1501                                         PointerType<ConstType<Size>, PrivateAS>,
1502                                         PointerType<ConstType<Size>, PrivateAS>]>;
1503 }
1504
1505
1506 //--------------------------------------------------------------------
1507 // End of the builtin functions defined in the OpenCL C specification.
1508 // Builtin functions defined in the OpenCL C Extension are below.
1509 //--------------------------------------------------------------------
1510
1511
1512 // OpenCL Extension v2.0 s9.18 - Mipmaps
1513 let Extension = FuncExtKhrMipmapImage in {
1514   // Added to section 6.13.14.2.
1515   foreach aQual = ["RO"] in {
1516     foreach imgTy = [Image2d] in {
1517       foreach name = ["read_imagef"] in {
1518         def : Builtin<name, [VectorType<Float, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, Float], Attr.Pure>;
1519         def : Builtin<name, [VectorType<Float, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, VectorType<Float, 2>, VectorType<Float, 2>], Attr.Pure>;
1520       }
1521       foreach name = ["read_imagei"] in {
1522         def : Builtin<name, [VectorType<Int, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, Float], Attr.Pure>;
1523         def : Builtin<name, [VectorType<Int, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, VectorType<Float, 2>, VectorType<Float, 2>], Attr.Pure>;
1524       }
1525       foreach name = ["read_imageui"] in {
1526         def : Builtin<name, [VectorType<UInt, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, Float], Attr.Pure>;
1527         def : Builtin<name, [VectorType<UInt, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, VectorType<Float, 2>, VectorType<Float, 2>], Attr.Pure>;
1528       }
1529     }
1530     foreach imgTy = [Image2dDepth] in {
1531       foreach name = ["read_imagef"] in {
1532         def : Builtin<name, [Float, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, Float], Attr.Pure>;
1533         def : Builtin<name, [Float, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, VectorType<Float, 2>, VectorType<Float, 2>], Attr.Pure>;
1534       }
1535     }
1536     foreach imgTy = [Image1d] in {
1537       foreach name = ["read_imagef"] in {
1538         def : Builtin<name, [VectorType<Float, 4>, ImageType<imgTy, aQual>, Sampler, Float, Float], Attr.Pure>;
1539         def : Builtin<name, [VectorType<Float, 4>, ImageType<imgTy, aQual>, Sampler, Float, Float, Float], Attr.Pure>;
1540       }
1541       foreach name = ["read_imagei"] in {
1542         def : Builtin<name, [VectorType<Int, 4>, ImageType<imgTy, aQual>, Sampler, Float, Float], Attr.Pure>;
1543         def : Builtin<name, [VectorType<Int, 4>, ImageType<imgTy, aQual>, Sampler, Float, Float, Float], Attr.Pure>;
1544       }
1545       foreach name = ["read_imageui"] in {
1546         def : Builtin<name, [VectorType<UInt, 4>, ImageType<imgTy, aQual>, Sampler, Float, Float], Attr.Pure>;
1547         def : Builtin<name, [VectorType<UInt, 4>, ImageType<imgTy, aQual>, Sampler, Float, Float, Float], Attr.Pure>;
1548       }
1549     }
1550     foreach imgTy = [Image3d] in {
1551       foreach name = ["read_imagef"] in {
1552         def : Builtin<name, [VectorType<Float, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, VectorType<Float, 4>, VectorType<Float, 4>], Attr.Pure>;
1553         def : Builtin<name, [VectorType<Float, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, Float], Attr.Pure>;
1554       }
1555       foreach name = ["read_imagei"] in {
1556         def : Builtin<name, [VectorType<Int, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, VectorType<Float, 4>, VectorType<Float, 4>], Attr.Pure>;
1557         def : Builtin<name, [VectorType<Int, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, Float], Attr.Pure>;
1558       }
1559       foreach name = ["read_imageui"] in {
1560         def : Builtin<name, [VectorType<UInt, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, VectorType<Float, 4>, VectorType<Float, 4>], Attr.Pure>;
1561         def : Builtin<name, [VectorType<UInt, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, Float], Attr.Pure>;
1562       }
1563     }
1564     foreach imgTy = [Image1dArray] in {
1565       foreach name = ["read_imagef"] in {
1566         def : Builtin<name, [VectorType<Float, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, Float], Attr.Pure>;
1567         def : Builtin<name, [VectorType<Float, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, Float, Float], Attr.Pure>;
1568       }
1569       foreach name = ["read_imagei"] in {
1570         def : Builtin<name, [VectorType<Int, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, Float], Attr.Pure>;
1571         def : Builtin<name, [VectorType<Int, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, Float, Float], Attr.Pure>;
1572       }
1573       foreach name = ["read_imageui"] in {
1574         def : Builtin<name, [VectorType<UInt, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, Float], Attr.Pure>;
1575         def : Builtin<name, [VectorType<UInt, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 2>, Float, Float], Attr.Pure>;
1576       }
1577     }
1578     foreach imgTy = [Image2dArray] in {
1579       foreach name = ["read_imagef"] in {
1580         def : Builtin<name, [VectorType<Float, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, Float], Attr.Pure>;
1581         def : Builtin<name, [VectorType<Float, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, VectorType<Float, 2>, VectorType<Float, 2>], Attr.Pure>;
1582       }
1583       foreach name = ["read_imagei"] in {
1584         def : Builtin<name, [VectorType<Int, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, Float], Attr.Pure>;
1585         def : Builtin<name, [VectorType<Int, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, VectorType<Float, 2>, VectorType<Float, 2>], Attr.Pure>;
1586       }
1587       foreach name = ["read_imageui"] in {
1588         def : Builtin<name, [VectorType<UInt, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, Float], Attr.Pure>;
1589         def : Builtin<name, [VectorType<UInt, 4>, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, VectorType<Float, 2>, VectorType<Float, 2>], Attr.Pure>;
1590       }
1591     }
1592     foreach imgTy = [Image2dArrayDepth] in {
1593       foreach name = ["read_imagef"] in {
1594         def : Builtin<name, [Float, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, Float], Attr.Pure>;
1595         def : Builtin<name, [Float, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, VectorType<Float, 2>, VectorType<Float, 2>], Attr.Pure>;
1596       }
1597     }
1598   }
1599 }
1600
1601 // Added to section 6.13.14.5
1602 multiclass ImageQueryNumMipLevels<string aQual> {
1603   foreach imgTy = [Image1d, Image2d, Image3d, Image1dArray, Image2dArray, Image2dDepth, Image2dArrayDepth] in {
1604     def : Builtin<"get_image_num_mip_levels", [Int, ImageType<imgTy, aQual>]>;
1605   }
1606 }
1607
1608 let Extension = FuncExtKhrMipmapImage in {
1609   defm : ImageQueryNumMipLevels<"RO">;
1610   defm : ImageQueryNumMipLevels<"WO">;
1611   defm : ImageQueryNumMipLevels<"RW">;
1612 }
1613
1614 // Write functions are enabled using a separate extension.
1615 let Extension = FuncExtKhrMipmapImageWrites in {
1616   // Added to section 6.13.14.4.
1617   foreach aQual = ["WO"] in {
1618     foreach imgTy = [Image2d] in {
1619       def : Builtin<"write_imagef", [Void, ImageType<imgTy, aQual>, VectorType<Int, 2>, Int, VectorType<Float, 4>]>;
1620       def : Builtin<"write_imagei", [Void, ImageType<imgTy, aQual>, VectorType<Int, 2>, Int, VectorType<Int, 4>]>;
1621       def : Builtin<"write_imageui", [Void, ImageType<imgTy, aQual>, VectorType<Int, 2>, Int, VectorType<UInt, 4>]>;
1622     }
1623     def : Builtin<"write_imagef", [Void, ImageType<Image2dDepth, aQual>, VectorType<Int, 2>, Int, Float]>;
1624     foreach imgTy = [Image1d] in {
1625       def : Builtin<"write_imagef", [Void, ImageType<imgTy, aQual>, Int, Int, VectorType<Float, 4>]>;
1626       def : Builtin<"write_imagei", [Void, ImageType<imgTy, aQual>, Int, Int, VectorType<Int, 4>]>;
1627       def : Builtin<"write_imageui", [Void, ImageType<imgTy, aQual>, Int, Int, VectorType<UInt, 4>]>;
1628     }
1629     foreach imgTy = [Image1dArray] in {
1630       def : Builtin<"write_imagef", [Void, ImageType<imgTy, aQual>, VectorType<Int, 2>, Int, VectorType<Float, 4>]>;
1631       def : Builtin<"write_imagei", [Void, ImageType<imgTy, aQual>, VectorType<Int, 2>, Int, VectorType<Int, 4>]>;
1632       def : Builtin<"write_imageui", [Void, ImageType<imgTy, aQual>, VectorType<Int, 2>, Int, VectorType<UInt, 4>]>;
1633     }
1634     foreach imgTy = [Image2dArray] in {
1635       def : Builtin<"write_imagef", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int, VectorType<Float, 4>]>;
1636       def : Builtin<"write_imagei", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int, VectorType<Int, 4>]>;
1637       def : Builtin<"write_imageui", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int, VectorType<UInt, 4>]>;
1638     }
1639     def : Builtin<"write_imagef", [Void, ImageType<Image2dArrayDepth, aQual>, VectorType<Int, 4>, Int, Float]>;
1640     foreach imgTy = [Image3d] in {
1641       def : Builtin<"write_imagef", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int, VectorType<Float, 4>]>;
1642       def : Builtin<"write_imagei", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int, VectorType<Int, 4>]>;
1643       def : Builtin<"write_imageui", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int, VectorType<UInt, 4>]>;
1644     }
1645   }
1646 }
1647
1648 //--------------------------------------------------------------------
1649 // OpenCL Extension v2.0 s18.3 - Creating OpenCL Memory Objects from OpenGL MSAA Textures
1650 // --- Table 6.13.14.3 ---
1651 multiclass ImageReadMsaa<string aQual> {
1652   foreach imgTy = [Image2dMsaa] in {
1653     def : Builtin<"read_imagef", [VectorType<Float, 4>, ImageType<imgTy, aQual>, VectorType<Int, 2>, Int], Attr.Pure>;
1654     def : Builtin<"read_imagei", [VectorType<Int, 4>, ImageType<imgTy, aQual>, VectorType<Int, 2>, Int], Attr.Pure>;
1655     def : Builtin<"read_imageui", [VectorType<UInt, 4>, ImageType<imgTy, aQual>, VectorType<Int, 2>, Int], Attr.Pure>;
1656   }
1657   foreach imgTy = [Image2dArrayMsaa] in {
1658     def : Builtin<"read_imagef", [VectorType<Float, 4>, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int], Attr.Pure>;
1659     def : Builtin<"read_imagei", [VectorType<Int, 4>, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int], Attr.Pure>;
1660     def : Builtin<"read_imageui", [VectorType<UInt, 4>, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int], Attr.Pure>;
1661   }
1662   foreach name = ["read_imagef"] in {
1663     def : Builtin<name, [Float, ImageType<Image2dMsaaDepth, aQual>, VectorType<Int, 2>, Int], Attr.Pure>;
1664     def : Builtin<name, [Float, ImageType<Image2dArrayMsaaDepth, aQual>, VectorType<Int, 4>, Int], Attr.Pure>;
1665   }
1666 }
1667
1668 // --- Table 6.13.14.5 ---
1669 multiclass ImageQueryMsaa<string aQual> {
1670   foreach imgTy = [Image2dMsaa, Image2dArrayMsaa, Image2dMsaaDepth, Image2dArrayMsaaDepth] in {
1671     foreach name = ["get_image_width", "get_image_height",
1672                     "get_image_channel_data_type", "get_image_channel_order",
1673                     "get_image_num_samples"] in {
1674       def : Builtin<name, [Int, ImageType<imgTy, aQual>], Attr.Const>;
1675     }
1676     def : Builtin<"get_image_dim", [VectorType<Int, 2>, ImageType<imgTy, aQual>], Attr.Const>;
1677   }
1678   foreach imgTy = [Image2dArrayMsaa, Image2dArrayMsaaDepth] in {
1679     def : Builtin<"get_image_array_size", [Size, ImageType<imgTy, aQual>], Attr.Const>;
1680   }
1681 }
1682
1683 let Extension = FuncExtKhrGlMsaaSharing in {
1684   defm : ImageReadMsaa<"RO">;
1685   defm : ImageQueryMsaa<"RO">;
1686   defm : ImageQueryMsaa<"WO">;
1687   defm : ImageReadMsaa<"RW">;
1688   defm : ImageQueryMsaa<"RW">;
1689 }
1690
1691 //--------------------------------------------------------------------
1692 // OpenCL Extension v2.0 s28 - Subgroups
1693 // --- Table 28.2.1 ---
1694 let Extension = FuncExtKhrSubgroups in {
1695   foreach name = ["get_sub_group_size", "get_max_sub_group_size",
1696                   "get_num_sub_groups", "get_sub_group_id",
1697                   "get_sub_group_local_id"] in {
1698     def : Builtin<name, [UInt]>;
1699   }
1700   let MinVersion = CL20 in {
1701     foreach name = ["get_enqueued_num_sub_groups"] in {
1702       def : Builtin<name, [UInt]>;
1703     }
1704   }
1705 }
1706
1707 // --- Table 28.2.2 ---
1708 let Extension = FuncExtKhrSubgroups in {
1709   def : Builtin<"sub_group_barrier", [Void, MemFenceFlags], Attr.Convergent>;
1710   let MinVersion = CL20 in {
1711     def : Builtin<"sub_group_barrier", [Void, MemFenceFlags, MemoryScope], Attr.Convergent>;
1712   }
1713 }
1714
1715 // --- Table 28.2.4 ---
1716 let Extension = FuncExtKhrSubgroups in {
1717   foreach name = ["sub_group_all", "sub_group_any"] in {
1718     def : Builtin<name, [Int, Int], Attr.Convergent>;
1719   }
1720   foreach name = ["sub_group_broadcast"] in {
1721     def : Builtin<name, [IntLongFloatGenType1, IntLongFloatGenType1, UInt], Attr.Convergent>;
1722   }
1723   foreach name = ["sub_group_reduce_", "sub_group_scan_exclusive_",
1724                   "sub_group_scan_inclusive_"] in {
1725     foreach op = ["add", "min", "max"] in {
1726       def : Builtin<name # op, [IntLongFloatGenType1, IntLongFloatGenType1], Attr.Convergent>;
1727     }
1728   }
1729 }
1730
1731 // OpenCL Extension v3.0 s38 - Extended Subgroup Functions
1732
1733 // Section 38.4.1 - cl_khr_subgroup_extended_types
1734 let Extension = FuncExtKhrSubgroupExtendedTypes in {
1735   // For sub_group_broadcast, add scalar char, uchar, short, and ushort support,
1736   def : Builtin<"sub_group_broadcast", [CharShortGenType1, CharShortGenType1, UInt], Attr.Convergent>;
1737   // gentype may additionally be one of the supported built-in vector data types.
1738   def : Builtin<"sub_group_broadcast", [AGenTypeNNoScalar, AGenTypeNNoScalar, UInt], Attr.Convergent>;
1739
1740   foreach name = ["sub_group_reduce_", "sub_group_scan_exclusive_",
1741                   "sub_group_scan_inclusive_"] in {
1742     foreach op = ["add", "min", "max"] in {
1743       def : Builtin<name # op, [CharShortGenType1, CharShortGenType1], Attr.Convergent>;
1744     }
1745   }
1746 }
1747
1748 // Section 38.5.1 - cl_khr_subgroup_non_uniform_vote
1749 let Extension = FuncExtKhrSubgroupNonUniformVote in {
1750   def : Builtin<"sub_group_elect", [Int]>;
1751   def : Builtin<"sub_group_non_uniform_all", [Int, Int]>;
1752   def : Builtin<"sub_group_non_uniform_any", [Int, Int]>;
1753   def : Builtin<"sub_group_non_uniform_all_equal", [Int, AGenType1]>;
1754 }
1755
1756 // Section 38.6.1 - cl_khr_subgroup_ballot
1757 let Extension = FuncExtKhrSubgroupBallot in {
1758   def : Builtin<"sub_group_non_uniform_broadcast", [AGenTypeN, AGenTypeN, UInt]>;
1759   def : Builtin<"sub_group_broadcast_first", [AGenType1, AGenType1]>;
1760   def : Builtin<"sub_group_ballot", [VectorType<UInt, 4>, Int]>;
1761   def : Builtin<"sub_group_inverse_ballot", [Int, VectorType<UInt, 4>], Attr.Const>;
1762   def : Builtin<"sub_group_ballot_bit_extract", [Int, VectorType<UInt, 4>, UInt], Attr.Const>;
1763   def : Builtin<"sub_group_ballot_bit_count", [UInt, VectorType<UInt, 4>], Attr.Const>;
1764   def : Builtin<"sub_group_ballot_inclusive_scan", [UInt, VectorType<UInt, 4>]>;
1765   def : Builtin<"sub_group_ballot_exclusive_scan", [UInt, VectorType<UInt, 4>]>;
1766   def : Builtin<"sub_group_ballot_find_lsb", [UInt, VectorType<UInt, 4>]>;
1767   def : Builtin<"sub_group_ballot_find_msb", [UInt, VectorType<UInt, 4>]>;
1768
1769   foreach op = ["eq", "ge", "gt", "le", "lt"] in {
1770     def : Builtin<"get_sub_group_" # op # "_mask", [VectorType<UInt, 4>], Attr.Const>;
1771   }
1772 }
1773
1774 // Section 38.7.1 - cl_khr_subgroup_non_uniform_arithmetic
1775 let Extension = FuncExtKhrSubgroupNonUniformArithmetic in {
1776   foreach name = ["reduce_", "scan_exclusive_", "scan_inclusive_"] in {
1777     foreach op = ["add", "min", "max", "mul"] in {
1778       def : Builtin<"sub_group_non_uniform_" # name # op, [AGenType1, AGenType1]>;
1779     }
1780     foreach op = ["and", "or", "xor"] in {
1781       def : Builtin<"sub_group_non_uniform_" # name # op, [AIGenType1, AIGenType1]>;
1782     }
1783     foreach op = ["and", "or", "xor"] in {
1784       def : Builtin<"sub_group_non_uniform_" # name # "logical_" # op, [Int, Int]>;
1785     }
1786   }
1787 }
1788
1789 // Section 38.8.1 - cl_khr_subgroup_shuffle
1790 let Extension = FuncExtKhrSubgroupShuffle in {
1791   def : Builtin<"sub_group_shuffle", [AGenType1, AGenType1, UInt]>;
1792   def : Builtin<"sub_group_shuffle_xor", [AGenType1, AGenType1, UInt]>;
1793 }
1794
1795 // Section 38.9.1 - cl_khr_subgroup_shuffle_relative
1796 let Extension = FuncExtKhrSubgroupShuffleRelative in {
1797   def : Builtin<"sub_group_shuffle_up", [AGenType1, AGenType1, UInt]>;
1798   def : Builtin<"sub_group_shuffle_down", [AGenType1, AGenType1, UInt]>;
1799 }
1800
1801 // Section 38.10.1 - cl_khr_subgroup_clustered_reduce
1802 let Extension = FuncExtKhrSubgroupClusteredReduce in {
1803   foreach op = ["add", "min", "max", "mul"] in {
1804     def : Builtin<"sub_group_clustered_reduce_" # op, [AGenType1, AGenType1, UInt]>;
1805   }
1806   foreach op = ["and", "or", "xor"] in {
1807     def : Builtin<"sub_group_clustered_reduce_" # op, [AIGenType1, AIGenType1, UInt]>;
1808   }
1809   foreach op = ["and", "or", "xor"] in {
1810     def : Builtin<"sub_group_clustered_reduce_logical_" # op, [Int, Int, UInt]>;
1811   }
1812 }
1813
1814 // Section 40.3.1 - cl_khr_extended_bit_ops
1815 let Extension = FuncExtKhrExtendedBitOps in {
1816   def : Builtin<"bitfield_insert", [AIGenTypeN, AIGenTypeN, AIGenTypeN, UInt, UInt], Attr.Const>;
1817   def : Builtin<"bitfield_extract_signed", [SGenTypeN, SGenTypeN, UInt, UInt], Attr.Const>;
1818   def : Builtin<"bitfield_extract_signed", [SGenTypeN, UGenTypeN, UInt, UInt], Attr.Const>;
1819   def : Builtin<"bitfield_extract_unsigned", [UGenTypeN, SGenTypeN, UInt, UInt], Attr.Const>;
1820   def : Builtin<"bitfield_extract_unsigned", [UGenTypeN, UGenTypeN, UInt, UInt], Attr.Const>;
1821   def : Builtin<"bit_reverse", [AIGenTypeN, AIGenTypeN], Attr.Const>;
1822 }
1823
1824 // Section 42.3 - cl_khr_integer_dot_product
1825 let Extension = FunctionExtension<"__opencl_c_integer_dot_product_input_4x8bit"> in {
1826   def : Builtin<"dot", [UInt, VectorType<UChar, 4>, VectorType<UChar, 4>], Attr.Const>;
1827   def : Builtin<"dot", [Int, VectorType<Char, 4>, VectorType<Char, 4>], Attr.Const>;
1828   def : Builtin<"dot", [Int, VectorType<UChar, 4>, VectorType<Char, 4>], Attr.Const>;
1829   def : Builtin<"dot", [Int, VectorType<Char, 4>, VectorType<UChar, 4>], Attr.Const>;
1830
1831   def : Builtin<"dot_acc_sat", [UInt, VectorType<UChar, 4>, VectorType<UChar, 4>, UInt], Attr.Const>;
1832   def : Builtin<"dot_acc_sat", [Int, VectorType<Char, 4>, VectorType<Char, 4>, Int], Attr.Const>;
1833   def : Builtin<"dot_acc_sat", [Int, VectorType<UChar, 4>, VectorType<Char, 4>, Int], Attr.Const>;
1834   def : Builtin<"dot_acc_sat", [Int, VectorType<Char, 4>, VectorType<UChar, 4>, Int], Attr.Const>;
1835 }
1836
1837 let Extension = FunctionExtension<"__opencl_c_integer_dot_product_input_4x8bit_packed"> in {
1838   def : Builtin<"dot_4x8packed_uu_uint", [UInt, UInt, UInt], Attr.Const>;
1839   def : Builtin<"dot_4x8packed_ss_int", [Int, UInt, UInt], Attr.Const>;
1840   def : Builtin<"dot_4x8packed_us_int", [Int, UInt, UInt], Attr.Const>;
1841   def : Builtin<"dot_4x8packed_su_int", [Int, UInt, UInt], Attr.Const>;
1842
1843   def : Builtin<"dot_acc_sat_4x8packed_uu_uint", [UInt, UInt, UInt, UInt], Attr.Const>;
1844   def : Builtin<"dot_acc_sat_4x8packed_ss_int", [Int, UInt, UInt, Int], Attr.Const>;
1845   def : Builtin<"dot_acc_sat_4x8packed_us_int", [Int, UInt, UInt, Int], Attr.Const>;
1846   def : Builtin<"dot_acc_sat_4x8packed_su_int", [Int, UInt, UInt, Int], Attr.Const>;
1847 }
1848
1849 // Section 48.3 - cl_khr_subgroup_rotate
1850 let Extension = FunctionExtension<"cl_khr_subgroup_rotate"> in {
1851   def : Builtin<"sub_group_rotate", [AGenType1, AGenType1, Int], Attr.Convergent>;
1852   def : Builtin<"sub_group_clustered_rotate", [AGenType1, AGenType1, Int, UInt], Attr.Convergent>;
1853 }
1854
1855 //--------------------------------------------------------------------
1856 // Arm extensions.
1857 let Extension = ArmIntegerDotProductInt8 in {
1858   foreach name = ["arm_dot"] in {
1859     def : Builtin<name, [UInt, VectorType<UChar, 4>, VectorType<UChar, 4>]>;
1860     def : Builtin<name, [Int, VectorType<Char, 4>, VectorType<Char, 4>]>;
1861   }
1862 }
1863 let Extension = ArmIntegerDotProductAccumulateInt8 in {
1864   foreach name = ["arm_dot_acc"] in {
1865     def : Builtin<name, [UInt, VectorType<UChar, 4>, VectorType<UChar, 4>, UInt]>;
1866     def : Builtin<name, [Int, VectorType<Char, 4>, VectorType<Char, 4>, Int]>;
1867   }
1868 }
1869 let Extension = ArmIntegerDotProductAccumulateInt16 in {
1870   foreach name = ["arm_dot_acc"] in {
1871     def : Builtin<name, [UInt, VectorType<UShort, 2>, VectorType<UShort, 2>, UInt]>;
1872     def : Builtin<name, [Int, VectorType<Short, 2>, VectorType<Short, 2>, Int]>;
1873   }
1874 }
1875 let Extension = ArmIntegerDotProductAccumulateSaturateInt8 in {
1876   foreach name = ["arm_dot_acc_sat"] in {
1877     def : Builtin<name, [UInt, VectorType<UChar, 4>, VectorType<UChar, 4>, UInt]>;
1878     def : Builtin<name, [Int, VectorType<Char, 4>, VectorType<Char, 4>, Int]>;
1879   }
1880 }