CbC/CbC_llvm: clang/lib/Sema/OpenCLBuiltins.td comparison

comparison clang/lib/Sema/OpenCLBuiltins.td @ 236:c4bab56944e8 llvm-original

LLVM 16

author	kono
date	Wed, 09 Nov 2022 17:45:10 +0900
parents	79ff65ed7e25
children

comparison

equal deleted inserted replaced

-:70dce7da266c
+:c4bab56944e8
 // During overload resolution, when a builtin function overload contains a type
 // with a TypeExtension, those overloads are skipped when the extension is
 // disabled.
 class TypeExtension<string _Ext> : AbstractExtension<_Ext>;
+// Concatenate zero or more space-separated extensions in NewExts to Base and
+// return the resulting FunctionExtension in ret.
+class concatExtension<FunctionExtension Base, string NewExts> {
+FunctionExtension ret = FunctionExtension<
+!cond(
+// Return Base extension if NewExts is empty,
+!empty(NewExts) : Base.ExtName,
+// otherwise, return NewExts if Base extension is empty,
+!empty(Base.ExtName) : NewExts,
+// otherwise, concatenate NewExts to Base.
+true : Base.ExtName # " " # NewExts
+)
+>;
+}
 // TypeExtension definitions.
 def NoTypeExt   : TypeExtension<"">;
 def Fp16TypeExt : TypeExtension<"cl_khr_fp16">;
 def Fp64TypeExt : TypeExtension<"cl_khr_fp64">;
+def Atomic64TypeExt : TypeExtension<"cl_khr_int64_base_atomics cl_khr_int64_extended_atomics">;
+def AtomicFp64TypeExt : TypeExtension<"cl_khr_int64_base_atomics cl_khr_int64_extended_atomics cl_khr_fp64">;
 // FunctionExtension definitions.
 def FuncExtNone                          : FunctionExtension<"">;
-def FuncExtKhrSubgroups                  : FunctionExtension<"cl_khr_subgroups">;
+def FuncExtKhrSubgroups                  : FunctionExtension<"__opencl_subgroup_builtins">;
 def FuncExtKhrSubgroupExtendedTypes      : FunctionExtension<"cl_khr_subgroup_extended_types">;
 def FuncExtKhrSubgroupNonUniformVote     : FunctionExtension<"cl_khr_subgroup_non_uniform_vote">;
 def FuncExtKhrSubgroupBallot             : FunctionExtension<"cl_khr_subgroup_ballot">;
 def FuncExtKhrSubgroupNonUniformArithmetic: FunctionExtension<"cl_khr_subgroup_non_uniform_arithmetic">;
 def FuncExtKhrSubgroupShuffle            : FunctionExtension<"cl_khr_subgroup_shuffle">;
 def FuncExtKhrSubgroupShuffleRelative    : FunctionExtension<"cl_khr_subgroup_shuffle_relative">;
 def FuncExtKhrSubgroupClusteredReduce    : FunctionExtension<"cl_khr_subgroup_clustered_reduce">;
+def FuncExtKhrExtendedBitOps             : FunctionExtension<"cl_khr_extended_bit_ops">;
 def FuncExtKhrGlobalInt32BaseAtomics     : FunctionExtension<"cl_khr_global_int32_base_atomics">;
 def FuncExtKhrGlobalInt32ExtendedAtomics : FunctionExtension<"cl_khr_global_int32_extended_atomics">;
 def FuncExtKhrLocalInt32BaseAtomics      : FunctionExtension<"cl_khr_local_int32_base_atomics">;
 def FuncExtKhrLocalInt32ExtendedAtomics  : FunctionExtension<"cl_khr_local_int32_extended_atomics">;
 def FuncExtKhrInt64BaseAtomics           : FunctionExtension<"cl_khr_int64_base_atomics">;
 def FuncExtKhrInt64ExtendedAtomics       : FunctionExtension<"cl_khr_int64_extended_atomics">;
 def FuncExtKhrMipmapImage                : FunctionExtension<"cl_khr_mipmap_image">;
 def FuncExtKhrMipmapImageWrites          : FunctionExtension<"cl_khr_mipmap_image_writes">;
 def FuncExtKhrGlMsaaSharing              : FunctionExtension<"cl_khr_gl_msaa_sharing">;
+def FuncExtOpenCLCDeviceEnqueue          : FunctionExtension<"__opencl_c_device_enqueue">;
+def FuncExtOpenCLCGenericAddressSpace    : FunctionExtension<"__opencl_c_generic_address_space">;
+def FuncExtOpenCLCNamedAddressSpaceBuiltins : FunctionExtension<"__opencl_c_named_address_space_builtins">;
+def FuncExtOpenCLCPipes                  : FunctionExtension<"__opencl_c_pipes">;
+def FuncExtOpenCLCWGCollectiveFunctions  : FunctionExtension<"__opencl_c_work_group_collective_functions">;
+def FuncExtOpenCLCReadWriteImages        : FunctionExtension<"__opencl_c_read_write_images">;
+def FuncExtFloatAtomicsFp16GlobalASLoadStore  : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_global_atomic_load_store">;
+def FuncExtFloatAtomicsFp16LocalASLoadStore   : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_local_atomic_load_store">;
+def FuncExtFloatAtomicsFp16GenericASLoadStore : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_global_atomic_load_store __opencl_c_ext_fp16_local_atomic_load_store">;
+def FuncExtFloatAtomicsFp16GlobalASAdd        : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_global_atomic_add">;
+def FuncExtFloatAtomicsFp32GlobalASAdd        : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp32_global_atomic_add">;
+def FuncExtFloatAtomicsFp64GlobalASAdd        : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp64_global_atomic_add">;
+def FuncExtFloatAtomicsFp16LocalASAdd         : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_local_atomic_add">;
+def FuncExtFloatAtomicsFp32LocalASAdd         : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp32_local_atomic_add">;
+def FuncExtFloatAtomicsFp64LocalASAdd         : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp64_local_atomic_add">;
+def FuncExtFloatAtomicsFp16GenericASAdd       : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_local_atomic_add __opencl_c_ext_fp16_global_atomic_add">;
+def FuncExtFloatAtomicsFp32GenericASAdd       : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp32_local_atomic_add __opencl_c_ext_fp32_global_atomic_add">;
+def FuncExtFloatAtomicsFp64GenericASAdd       : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp64_local_atomic_add __opencl_c_ext_fp64_global_atomic_add">;
+def FuncExtFloatAtomicsFp16GlobalASMinMax     : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_global_atomic_min_max">;
+def FuncExtFloatAtomicsFp32GlobalASMinMax     : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp32_global_atomic_min_max">;
+def FuncExtFloatAtomicsFp64GlobalASMinMax     : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp64_global_atomic_min_max">;
+def FuncExtFloatAtomicsFp16LocalASMinMax      : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_local_atomic_min_max">;
+def FuncExtFloatAtomicsFp32LocalASMinMax      : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp32_local_atomic_min_max">;
+def FuncExtFloatAtomicsFp64LocalASMinMax      : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp64_local_atomic_min_max">;
+def FuncExtFloatAtomicsFp16GenericASMinMax    : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_local_atomic_min_max __opencl_c_ext_fp16_global_atomic_min_max">;
+def FuncExtFloatAtomicsFp32GenericASMinMax    : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp32_local_atomic_min_max __opencl_c_ext_fp32_global_atomic_min_max">;
+def FuncExtFloatAtomicsFp64GenericASMinMax    : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp64_local_atomic_min_max __opencl_c_ext_fp64_global_atomic_min_max">;
 // Not a real extension, but a workaround to add C++ for OpenCL specific builtins.
 def FuncExtOpenCLCxx                     : FunctionExtension<"__cplusplus">;
-// Multiple extensions
-def FuncExtKhrMipmapWritesAndWrite3d     : FunctionExtension<"cl_khr_mipmap_image_writes cl_khr_3d_image_writes">;
 // Arm extensions.
 def ArmIntegerDotProductInt8                   : FunctionExtension<"cl_arm_integer_dot_product_int8">;
 def ArmIntegerDotProductAccumulateInt8         : FunctionExtension<"cl_arm_integer_dot_product_accumulate_int8">;
 def ArmIntegerDotProductAccumulateInt16        : FunctionExtension<"cl_arm_integer_dot_product_accumulate_int16">;
 // Inherited fields
 let IsPointer = _Ty.IsPointer;
 let IsConst = _Ty.IsConst;
 let IsVolatile = _Ty.IsVolatile;
 let AddrSpace = _Ty.AddrSpace;
-let Extension = _Ty.Extension;
+// Add TypeExtensions for writable "image3d_t" and "read_write" image types.
+let Extension = !cond(
+!and(!eq(_Ty.Name, "image3d_t"), !eq(_AccessQualifier, "WO")) : TypeExtension<"cl_khr_3d_image_writes">,
+!and(!eq(_Ty.Name, "image3d_t"), !eq(_AccessQualifier, "RW")) : TypeExtension<"cl_khr_3d_image_writes __opencl_c_read_write_images">,
+!or(!eq(_Ty.Name, "image2d_depth_t"), !eq(_Ty.Name, "image2d_array_depth_t")) : TypeExtension<"cl_khr_depth_images">,
+!eq(_AccessQualifier, "RW") : TypeExtension<"__opencl_c_read_write_images">,
+true : _Ty.Extension);
 }
 // OpenCL enum type (e.g. memory_scope).
 class EnumType<string _Name> :
 Type<_Name, QualType<"getOpenCLEnumType(S, \"" # _Name # "\")", 0>> {
 def ULong     : Type<"ulong",     QualType<"Context.UnsignedLongTy">>;
 def Float     : Type<"float",     QualType<"Context.FloatTy">>;
 let Extension = Fp64TypeExt in {
 def Double    : Type<"double",    QualType<"Context.DoubleTy">>;
 }
+// The half type for builtins that require the cl_khr_fp16 extension.
 let Extension = Fp16TypeExt in {
 def Half      : Type<"half",      QualType<"Context.HalfTy">>;
 }
+// Without the cl_khr_fp16 extension, the half type can only be used to declare
+// a pointer.  Define const and non-const pointer types in all address spaces.
+// Use the "__half" alias to allow the TableGen emitter to distinguish the
+// (extensionless) pointee type of these pointer-to-half types from the "half"
+// type defined above that already carries the cl_khr_fp16 extension.
+foreach AS = [PrivateAS, GlobalAS, ConstantAS, LocalAS, GenericAS] in {
+def "HalfPtr" # AS      : PointerType<Type<"__half", QualType<"Context.HalfTy">>, AS>;
+def "HalfPtrConst" # AS : PointerType<ConstType<Type<"__half", QualType<"Context.HalfTy">>>, AS>;
+}
 def Size      : Type<"size_t",    QualType<"Context.getSizeType()">>;
 def PtrDiff   : Type<"ptrdiff_t", QualType<"Context.getPointerDiffType()">>;
 def IntPtr    : Type<"intptr_t",  QualType<"Context.getIntPtrType()">>;
 def UIntPtr   : Type<"uintptr_t", QualType<"Context.getUIntPtrType()">>;
 def Void      : Type<"void",      QualType<"Context.VoidTy">>;
 def NDRange               : TypedefType<"ndrange_t">;
 // OpenCL v2.0 s6.13.11: Atomic integer and floating-point types.
 def AtomicInt             : Type<"atomic_int", QualType<"Context.getAtomicType(Context.IntTy)">>;
 def AtomicUInt            : Type<"atomic_uint", QualType<"Context.getAtomicType(Context.UnsignedIntTy)">>;
-def AtomicLong            : Type<"atomic_long", QualType<"Context.getAtomicType(Context.LongTy)">>;
+let Extension = Atomic64TypeExt in {
-def AtomicULong           : Type<"atomic_ulong", QualType<"Context.getAtomicType(Context.UnsignedLongTy)">>;
+def AtomicLong            : Type<"atomic_long", QualType<"Context.getAtomicType(Context.LongTy)">>;
+def AtomicULong           : Type<"atomic_ulong", QualType<"Context.getAtomicType(Context.UnsignedLongTy)">>;
+}
 def AtomicFloat           : Type<"atomic_float", QualType<"Context.getAtomicType(Context.FloatTy)">>;
-def AtomicDouble          : Type<"atomic_double", QualType<"Context.getAtomicType(Context.DoubleTy)">>;
+let Extension = AtomicFp64TypeExt in {
+def AtomicDouble          : Type<"atomic_double", QualType<"Context.getAtomicType(Context.DoubleTy)">>;
+}
+def AtomicHalf            : Type<"atomic_half", QualType<"Context.getAtomicType(Context.HalfTy)">>;
 def AtomicIntPtr          : Type<"atomic_intptr_t", QualType<"Context.getAtomicType(Context.getIntPtrType())">>;
 def AtomicUIntPtr         : Type<"atomic_uintptr_t", QualType<"Context.getAtomicType(Context.getUIntPtrType())">>;
 def AtomicSize            : Type<"atomic_size_t", QualType<"Context.getAtomicType(Context.getSizeType())">>;
 def AtomicPtrDiff         : Type<"atomic_ptrdiff_t", QualType<"Context.getAtomicType(Context.getPointerDiffType())">>;
 // --- 3 arguments ---
 foreach name = ["fma", "mad"] in {
 def : Builtin<name, [FGenTypeN, FGenTypeN, FGenTypeN, FGenTypeN], Attr.Const>;
 }
-// --- Version dependent ---
+// The following math builtins take pointer arguments.  Which overloads are
-let MaxVersion = CL20 in {
+// available depends on whether the generic address space feature is enabled.
-foreach AS = [GlobalAS, LocalAS, PrivateAS] in {
+multiclass MathWithPointer<list<AddressSpace> addrspaces> {
+foreach AS = addrspaces in {
 foreach name = ["fract", "modf", "sincos"] in {
 def : Builtin<name, [FGenTypeN, FGenTypeN, PointerType<FGenTypeN, AS>]>;
 }
 foreach name = ["frexp", "lgamma_r"] in {
 foreach Type = [GenTypeFloatVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeHalfVecAndScalar] in {
 def : Builtin<name, [Type, Type, Type, PointerType<GenTypeIntVecAndScalar, AS>]>;
 }
 }
 }
 }
-let MinVersion = CL20 in {
-foreach name = ["fract", "modf", "sincos"] in {
+let Extension = FuncExtOpenCLCNamedAddressSpaceBuiltins in {
-def : Builtin<name, [FGenTypeN, FGenTypeN, PointerType<FGenTypeN, GenericAS>]>;
+defm : MathWithPointer<[GlobalAS, LocalAS, PrivateAS]>;
 }
-foreach name = ["frexp", "lgamma_r"] in {
+let Extension = FuncExtOpenCLCGenericAddressSpace in {
-foreach Type = [GenTypeFloatVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeHalfVecAndScalar] in {
+defm : MathWithPointer<[GenericAS]>;
-def : Builtin<name, [Type, Type, PointerType<GenTypeIntVecAndScalar, GenericAS>]>;
-}  }
-foreach name = ["remquo"] in {
-foreach Type = [GenTypeFloatVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeHalfVecAndScalar] in {
-def : Builtin<name, [Type, Type, Type, PointerType<GenTypeIntVecAndScalar, GenericAS>]>;
-}
-}
 }
 // --- Table 9 ---
 foreach name = ["half_cos",
 "half_exp", "half_exp2", "half_exp10",
 //--------------------------------------------------------------------
 // OpenCL v1.1 s6.11.7, v1.2 s6.12.7, v2.0 s6.13.7 - Vector Data Load and Store Functions
 // 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
 // --- Table 15 ---
-// Variants for OpenCL versions below 2.0, using pointers to the global, local
+multiclass VloadVstore<list<AddressSpace> addrspaces, bit defStores> {
-// and private address spaces.
+foreach AS = addrspaces in {
-let MaxVersion = CL20 in {
-foreach AS = [GlobalAS, LocalAS, PrivateAS] in {
 foreach VSize = [2, 3, 4, 8, 16] in {
 foreach name = ["vload" # VSize] in {
-def : Builtin<name, [VectorType<Char, VSize>, Size, PointerType<ConstType<Char>, AS>]>;
+def : Builtin<name, [VectorType<Char, VSize>, Size, PointerType<ConstType<Char>, AS>], Attr.Pure>;
-def : Builtin<name, [VectorType<UChar, VSize>, Size, PointerType<ConstType<UChar>, AS>]>;
+def : Builtin<name, [VectorType<UChar, VSize>, Size, PointerType<ConstType<UChar>, AS>], Attr.Pure>;
-def : Builtin<name, [VectorType<Short, VSize>, Size, PointerType<ConstType<Short>, AS>]>;
+def : Builtin<name, [VectorType<Short, VSize>, Size, PointerType<ConstType<Short>, AS>], Attr.Pure>;
-def : Builtin<name, [VectorType<UShort, VSize>, Size, PointerType<ConstType<UShort>, AS>]>;
+def : Builtin<name, [VectorType<UShort, VSize>, Size, PointerType<ConstType<UShort>, AS>], Attr.Pure>;
-def : Builtin<name, [VectorType<Int, VSize>, Size, PointerType<ConstType<Int>, AS>]>;
+def : Builtin<name, [VectorType<Int, VSize>, Size, PointerType<ConstType<Int>, AS>], Attr.Pure>;
-def : Builtin<name, [VectorType<UInt, VSize>, Size, PointerType<ConstType<UInt>, AS>]>;
+def : Builtin<name, [VectorType<UInt, VSize>, Size, PointerType<ConstType<UInt>, AS>], Attr.Pure>;
-def : Builtin<name, [VectorType<Long, VSize>, Size, PointerType<ConstType<Long>, AS>]>;
+def : Builtin<name, [VectorType<Long, VSize>, Size, PointerType<ConstType<Long>, AS>], Attr.Pure>;
-def : Builtin<name, [VectorType<ULong, VSize>, Size, PointerType<ConstType<ULong>, AS>]>;
+def : Builtin<name, [VectorType<ULong, VSize>, Size, PointerType<ConstType<ULong>, AS>], Attr.Pure>;
-def : Builtin<name, [VectorType<Float, VSize>, Size, PointerType<ConstType<Float>, AS>]>;
+def : Builtin<name, [VectorType<Float, VSize>, Size, PointerType<ConstType<Float>, AS>], Attr.Pure>;
-def : Builtin<name, [VectorType<Double, VSize>, Size, PointerType<ConstType<Double>, AS>]>;
+def : Builtin<name, [VectorType<Double, VSize>, Size, PointerType<ConstType<Double>, AS>], Attr.Pure>;
-def : Builtin<name, [VectorType<Half, VSize>, Size, PointerType<ConstType<Half>, AS>]>;
+def : Builtin<name, [VectorType<Half, VSize>, Size, PointerType<ConstType<Half>, AS>], Attr.Pure>;
 }
-foreach name = ["vstore" # VSize] in {
+if defStores then {
-def : Builtin<name, [Void, VectorType<Char, VSize>, Size, PointerType<Char, AS>]>;
+foreach name = ["vstore" # VSize] in {
-def : Builtin<name, [Void, VectorType<UChar, VSize>, Size, PointerType<UChar, AS>]>;
+def : Builtin<name, [Void, VectorType<Char, VSize>, Size, PointerType<Char, AS>]>;
-def : Builtin<name, [Void, VectorType<Short, VSize>, Size, PointerType<Short, AS>]>;
+def : Builtin<name, [Void, VectorType<UChar, VSize>, Size, PointerType<UChar, AS>]>;
-def : Builtin<name, [Void, VectorType<UShort, VSize>, Size, PointerType<UShort, AS>]>;
+def : Builtin<name, [Void, VectorType<Short, VSize>, Size, PointerType<Short, AS>]>;
-def : Builtin<name, [Void, VectorType<Int, VSize>, Size, PointerType<Int, AS>]>;
+def : Builtin<name, [Void, VectorType<UShort, VSize>, Size, PointerType<UShort, AS>]>;
-def : Builtin<name, [Void, VectorType<UInt, VSize>, Size, PointerType<UInt, AS>]>;
+def : Builtin<name, [Void, VectorType<Int, VSize>, Size, PointerType<Int, AS>]>;
-def : Builtin<name, [Void, VectorType<Long, VSize>, Size, PointerType<Long, AS>]>;
+def : Builtin<name, [Void, VectorType<UInt, VSize>, Size, PointerType<UInt, AS>]>;
-def : Builtin<name, [Void, VectorType<ULong, VSize>, Size, PointerType<ULong, AS>]>;
+def : Builtin<name, [Void, VectorType<Long, VSize>, Size, PointerType<Long, AS>]>;
-def : Builtin<name, [Void, VectorType<Float, VSize>, Size, PointerType<Float, AS>]>;
+def : Builtin<name, [Void, VectorType<ULong, VSize>, Size, PointerType<ULong, AS>]>;
-def : Builtin<name, [Void, VectorType<Double, VSize>, Size, PointerType<Double, AS>]>;
+def : Builtin<name, [Void, VectorType<Float, VSize>, Size, PointerType<Float, AS>]>;
-def : Builtin<name, [Void, VectorType<Half, VSize>, Size, PointerType<Half, AS>]>;
+def : Builtin<name, [Void, VectorType<Double, VSize>, Size, PointerType<Double, AS>]>;
-}
+def : Builtin<name, [Void, VectorType<Half, VSize>, Size, PointerType<Half, AS>]>;
-foreach name = ["vloada_half" # VSize] in {
+}
-def : Builtin<name, [VectorType<Float, VSize>, Size, PointerType<ConstType<Half>, AS>]>;
+}
 }
+}
+}
+let Extension = FuncExtOpenCLCNamedAddressSpaceBuiltins in {
+defm : VloadVstore<[GlobalAS, LocalAS, PrivateAS], 1>;
+}
+let Extension = FuncExtOpenCLCGenericAddressSpace in {
+defm : VloadVstore<[GenericAS], 1>;
+}
+// vload with constant address space is available regardless of version.
+defm : VloadVstore<[ConstantAS], 0>;
+multiclass VloadVstoreHalf<list<AddressSpace> addrspaces, bit defStores> {
+foreach AS = addrspaces in {
+def : Builtin<"vload_half", [Float, Size, !cast<Type>("HalfPtrConst" # AS)], Attr.Pure>;
+foreach VSize = [2, 3, 4, 8, 16] in {
+foreach name = ["vload_half" # VSize, "vloada_half" # VSize] in {
+def : Builtin<name, [VectorType<Float, VSize>, Size, !cast<Type>("HalfPtrConst" # AS)], Attr.Pure>;
+}
+}
+if defStores then {
 foreach rnd = ["", "_rte", "_rtz", "_rtp", "_rtn"] in {
-foreach name = ["vstorea_half" # VSize # rnd] in {
+foreach name = ["vstore_half" # rnd] in {
-def : Builtin<name, [Void, VectorType<Float, VSize>, Size, PointerType<Half, AS>]>;
+def : Builtin<name, [Void, Float, Size, !cast<Type>("HalfPtr" # AS)]>;
-def : Builtin<name, [Void, VectorType<Double, VSize>, Size, PointerType<Half, AS>]>;
+def : Builtin<name, [Void, Double, Size, !cast<Type>("HalfPtr" # AS)]>;
 }
-}
+foreach VSize = [2, 3, 4, 8, 16] in {
-}
+foreach name = ["vstore_half" # VSize # rnd, "vstorea_half" # VSize # rnd] in {
-}
+def : Builtin<name, [Void, VectorType<Float, VSize>, Size, !cast<Type>("HalfPtr" # AS)]>;
-}
+def : Builtin<name, [Void, VectorType<Double, VSize>, Size, !cast<Type>("HalfPtr" # AS)]>;
-// Variants for OpenCL versions above 2.0, using pointers to the generic
+}
-// address space.
-let MinVersion = CL20 in {
-foreach VSize = [2, 3, 4, 8, 16] in {
-foreach name = ["vload" # VSize] in {
-def : Builtin<name, [VectorType<Char, VSize>, Size, PointerType<ConstType<Char>, GenericAS>]>;
-def : Builtin<name, [VectorType<UChar, VSize>, Size, PointerType<ConstType<UChar>, GenericAS>]>;
-def : Builtin<name, [VectorType<Short, VSize>, Size, PointerType<ConstType<Short>, GenericAS>]>;
-def : Builtin<name, [VectorType<UShort, VSize>, Size, PointerType<ConstType<UShort>, GenericAS>]>;
-def : Builtin<name, [VectorType<Int, VSize>, Size, PointerType<ConstType<Int>, GenericAS>]>;
-def : Builtin<name, [VectorType<UInt, VSize>, Size, PointerType<ConstType<UInt>, GenericAS>]>;
-def : Builtin<name, [VectorType<Long, VSize>, Size, PointerType<ConstType<Long>, GenericAS>]>;
-def : Builtin<name, [VectorType<ULong, VSize>, Size, PointerType<ConstType<ULong>, GenericAS>]>;
-def : Builtin<name, [VectorType<Float, VSize>, Size, PointerType<ConstType<Float>, GenericAS>]>;
-def : Builtin<name, [VectorType<Double, VSize>, Size, PointerType<ConstType<Double>, GenericAS>]>;
-def : Builtin<name, [VectorType<Half, VSize>, Size, PointerType<ConstType<Half>, GenericAS>]>;
-}
-foreach name = ["vstore" # VSize] in {
-def : Builtin<name, [Void, VectorType<Char, VSize>, Size, PointerType<Char, GenericAS>]>;
-def : Builtin<name, [Void, VectorType<UChar, VSize>, Size, PointerType<UChar, GenericAS>]>;
-def : Builtin<name, [Void, VectorType<Short, VSize>, Size, PointerType<Short, GenericAS>]>;
-def : Builtin<name, [Void, VectorType<UShort, VSize>, Size, PointerType<UShort, GenericAS>]>;
-def : Builtin<name, [Void, VectorType<Int, VSize>, Size, PointerType<Int, GenericAS>]>;
-def : Builtin<name, [Void, VectorType<UInt, VSize>, Size, PointerType<UInt, GenericAS>]>;
-def : Builtin<name, [Void, VectorType<Long, VSize>, Size, PointerType<Long, GenericAS>]>;
-def : Builtin<name, [Void, VectorType<ULong, VSize>, Size, PointerType<ULong, GenericAS>]>;
-def : Builtin<name, [Void, VectorType<Float, VSize>, Size, PointerType<Float, GenericAS>]>;
-def : Builtin<name, [Void, VectorType<Double, VSize>, Size, PointerType<Double, GenericAS>]>;
-def : Builtin<name, [Void, VectorType<Half, VSize>, Size, PointerType<Half, GenericAS>]>;
-}
-foreach name = ["vloada_half" # VSize] in {
-def : Builtin<name, [VectorType<Float, VSize>, Size, PointerType<ConstType<Half>, GenericAS>]>;
-}
-foreach rnd = ["", "_rte", "_rtz", "_rtp", "_rtn"] in {
-foreach name = ["vstorea_half" # VSize # rnd] in {
-def : Builtin<name, [Void, VectorType<Float, VSize>, Size, PointerType<Half, GenericAS>]>;
-def : Builtin<name, [Void, VectorType<Double, VSize>, Size, PointerType<Half, GenericAS>]>;
-}
-}
-}
-}
-// Variants using pointers to the constant address space.
-foreach VSize = [2, 3, 4, 8, 16] in {
-foreach name = ["vload" # VSize] in {
-def : Builtin<name, [VectorType<Char, VSize>, Size, PointerType<ConstType<Char>, ConstantAS>]>;
-def : Builtin<name, [VectorType<UChar, VSize>, Size, PointerType<ConstType<UChar>, ConstantAS>]>;
-def : Builtin<name, [VectorType<Short, VSize>, Size, PointerType<ConstType<Short>, ConstantAS>]>;
-def : Builtin<name, [VectorType<UShort, VSize>, Size, PointerType<ConstType<UShort>, ConstantAS>]>;
-def : Builtin<name, [VectorType<Int, VSize>, Size, PointerType<ConstType<Int>, ConstantAS>]>;
-def : Builtin<name, [VectorType<UInt, VSize>, Size, PointerType<ConstType<UInt>, ConstantAS>]>;
-def : Builtin<name, [VectorType<Long, VSize>, Size, PointerType<ConstType<Long>, ConstantAS>]>;
-def : Builtin<name, [VectorType<ULong, VSize>, Size, PointerType<ConstType<ULong>, ConstantAS>]>;
-def : Builtin<name, [VectorType<Float, VSize>, Size, PointerType<ConstType<Float>, ConstantAS>]>;
-def : Builtin<name, [VectorType<Double, VSize>, Size, PointerType<ConstType<Double>, ConstantAS>]>;
-def : Builtin<name, [VectorType<Half, VSize>, Size, PointerType<ConstType<Half>, ConstantAS>]>;
-}
-foreach name = ["vloada_half" # VSize] in {
-def : Builtin<name, [VectorType<Float, VSize>, Size, PointerType<ConstType<Half>, ConstantAS>]>;
-}
-}
-let MaxVersion = CL20 in {
-foreach AS = [GlobalAS, LocalAS, PrivateAS] in {
-def : Builtin<"vload_half", [Float, Size, PointerType<ConstType<Half>, AS>]>;
-def : Builtin<"vloada_half", [Float, Size, PointerType<ConstType<Half>, AS>]>;
-foreach VSize = [2, 3, 4, 8, 16] in {
-foreach name = ["vload_half" # VSize] in {
-def : Builtin<name, [VectorType<Float, VSize>, Size, PointerType<ConstType<Half>, AS>]>;
-}
-}
-foreach rnd = ["", "_rte", "_rtz", "_rtp", "_rtn"] in {
-foreach name = ["vstore_half" # rnd, "vstorea_half" # rnd] in {
-def : Builtin<name, [Void, Float, Size, PointerType<Half, AS>]>;
-def : Builtin<name, [Void, Double, Size, PointerType<Half, AS>]>;
-}
-foreach VSize = [2, 3, 4, 8, 16] in {
-foreach name = ["vstore_half" # VSize # rnd] in {
-def : Builtin<name, [Void, VectorType<Float, VSize>, Size, PointerType<Half, AS>]>;
-def : Builtin<name, [Void, VectorType<Double, VSize>, Size, PointerType<Half, AS>]>;
 }
 }
 }
 }
 }
-let MinVersion = CL20 in {
-foreach AS = [GenericAS] in {
+let Extension = FuncExtOpenCLCNamedAddressSpaceBuiltins in {
-def : Builtin<"vload_half", [Float, Size, PointerType<ConstType<Half>, AS>]>;
+defm : VloadVstoreHalf<[GlobalAS, LocalAS, PrivateAS], 1>;
-def : Builtin<"vloada_half", [Float, Size, PointerType<ConstType<Half>, AS>]>;
+}
-foreach VSize = [2, 3, 4, 8, 16] in {
+let Extension = FuncExtOpenCLCGenericAddressSpace in {
-foreach name = ["vload_half" # VSize] in {
+defm : VloadVstoreHalf<[GenericAS], 1>;
-def : Builtin<name, [VectorType<Float, VSize>, Size, PointerType<ConstType<Half>, AS>]>;
+}
-}
+// vload_half and vloada_half with constant address space are available regardless of version.
-}
+defm : VloadVstoreHalf<[ConstantAS], 0>;
-foreach rnd = ["", "_rte", "_rtz", "_rtp", "_rtn"] in {
-foreach name = ["vstore_half" # rnd, "vstorea_half" # rnd] in {
-def : Builtin<name, [Void, Float, Size, PointerType<Half, AS>]>;
-def : Builtin<name, [Void, Double, Size, PointerType<Half, AS>]>;
-}
-foreach VSize = [2, 3, 4, 8, 16] in {
-foreach name = ["vstore_half" # VSize # rnd] in {
-def : Builtin<name, [Void, VectorType<Float, VSize>, Size, PointerType<Half, AS>]>;
-def : Builtin<name, [Void, VectorType<Double, VSize>, Size, PointerType<Half, AS>]>;
-}
-}
-}
-}
-}
-foreach AS = [ConstantAS] in {
-def : Builtin<"vload_half", [Float, Size, PointerType<ConstType<Half>, AS>]>;
-def : Builtin<"vloada_half", [Float, Size, PointerType<ConstType<Half>, AS>]>;
-foreach VSize = [2, 3, 4, 8, 16] in {
-foreach name = ["vload_half" # VSize] in {
-def : Builtin<name, [VectorType<Float, VSize>, Size, PointerType<ConstType<Half>, AS>]>;
-}
-}
-}
 // OpenCL v3.0 s6.15.8 - Synchronization Functions.
 def : Builtin<"barrier", [Void, MemFenceFlags], Attr.Convergent>;
 let MinVersion = CL20 in {
 def : Builtin<"work_group_barrier", [Void, MemFenceFlags], Attr.Convergent>;
 def : Builtin<"write_mem_fence", [Void, MemFenceFlags]>;
 // OpenCL v3.0 s6.15.10 - Address Space Qualifier Functions.
 // to_global, to_local, to_private are declared in Builtins.def.
-let MinVersion = CL20 in {
+let Extension = FuncExtOpenCLCGenericAddressSpace in {
 // The OpenCL 3.0 specification defines these with a "gentype" argument indicating any builtin
 // type or user-defined type, which cannot be represented currently.  Hence we slightly diverge
 // by providing only the following overloads with a void pointer.
 def : Builtin<"get_fence", [MemFenceFlags, PointerType<Void, GenericAS>]>;
 def : Builtin<"get_fence", [MemFenceFlags, PointerType<ConstType<Void>, GenericAS>]>;
 }
 }
 }
 // OpenCL v2.0 s6.13.11 - Atomic Functions.
-let MinVersion = CL20 in {
-def : Builtin<"atomic_work_item_fence", [Void, MemFenceFlags, MemoryOrder, MemoryScope]>;
+// An atomic builtin with 2 additional _explicit variants.
+multiclass BuiltinAtomicExplicit<string Name, list<Type> Types, FunctionExtension BaseExt> {
+// Without explicit MemoryOrder or MemoryScope.
+let Extension = concatExtension<BaseExt, "__opencl_c_atomic_order_seq_cst __opencl_c_atomic_scope_device">.ret in {
+def : Builtin<Name, Types>;
+}
+// With an explicit MemoryOrder argument.
+let Extension = concatExtension<BaseExt, "__opencl_c_atomic_scope_device">.ret in {
+def : Builtin<Name # "_explicit", !listconcat(Types, [MemoryOrder])>;
+}
+// With explicit MemoryOrder and MemoryScope arguments.
+let Extension = BaseExt in {
+def : Builtin<Name # "_explicit", !listconcat(Types, [MemoryOrder, MemoryScope])>;
+}
+}
+// OpenCL 2.0 atomic functions that have a pointer argument in a given address space.
+multiclass OpenCL2Atomics<AddressSpace addrspace, FunctionExtension BaseExt> {
 foreach TypePair = [[AtomicInt, Int], [AtomicUInt, UInt],
 [AtomicLong, Long], [AtomicULong, ULong],
 [AtomicFloat, Float], [AtomicDouble, Double]] in {
-def : Builtin<"atomic_init",
+let Extension = BaseExt in {
-[Void, PointerType<VolatileType<TypePair[0]>, GenericAS>, TypePair[1]]>;
+def : Builtin<"atomic_init",
-def : Builtin<"atomic_store",
+[Void, PointerType<VolatileType<TypePair[0]>, addrspace>, TypePair[1]]>;
-[Void, PointerType<VolatileType<TypePair[0]>, GenericAS>, TypePair[1]]>;
+}
-def : Builtin<"atomic_store_explicit",
+defm : BuiltinAtomicExplicit<"atomic_store",
-[Void, PointerType<VolatileType<TypePair[0]>, GenericAS>, TypePair[1], MemoryOrder]>;
+[Void, PointerType<VolatileType<TypePair[0]>, addrspace>, TypePair[1]], BaseExt>;
-def : Builtin<"atomic_store_explicit",
+defm : BuiltinAtomicExplicit<"atomic_load",
-[Void, PointerType<VolatileType<TypePair[0]>, GenericAS>, TypePair[1], MemoryOrder, MemoryScope]>;
+[TypePair[1], PointerType<VolatileType<TypePair[0]>, addrspace>], BaseExt>;
-def : Builtin<"atomic_load",
+defm : BuiltinAtomicExplicit<"atomic_exchange",
-[TypePair[1], PointerType<VolatileType<TypePair[0]>, GenericAS>]>;
+[TypePair[1], PointerType<VolatileType<TypePair[0]>, addrspace>, TypePair[1]], BaseExt>;
-def : Builtin<"atomic_load_explicit",
-[TypePair[1], PointerType<VolatileType<TypePair[0]>, GenericAS>, MemoryOrder]>;
-def : Builtin<"atomic_load_explicit",
-[TypePair[1], PointerType<VolatileType<TypePair[0]>, GenericAS>, MemoryOrder, MemoryScope]>;
-def : Builtin<"atomic_exchange",
-[TypePair[1], PointerType<VolatileType<TypePair[0]>, GenericAS>, TypePair[1]]>;
-def : Builtin<"atomic_exchange_explicit",
-[TypePair[1], PointerType<VolatileType<TypePair[0]>, GenericAS>, TypePair[1], MemoryOrder]>;
-def : Builtin<"atomic_exchange_explicit",
-[TypePair[1], PointerType<VolatileType<TypePair[0]>, GenericAS>, TypePair[1], MemoryOrder, MemoryScope]>;
 foreach Variant = ["weak", "strong"] in {
-def : Builtin<"atomic_compare_exchange_" # Variant,
+foreach exp_ptr_addrspace = !cond(
-[Bool, PointerType<VolatileType<TypePair[0]>, GenericAS>,
+!eq(BaseExt, FuncExtOpenCLCGenericAddressSpace): [GenericAS],
-PointerType<TypePair[1], GenericAS>, TypePair[1]]>;
+!eq(BaseExt, FuncExtOpenCLCNamedAddressSpaceBuiltins): [GlobalAS, LocalAS, PrivateAS])
-def : Builtin<"atomic_compare_exchange_" # Variant # "_explicit",
+in {
-[Bool, PointerType<VolatileType<TypePair[0]>, GenericAS>,
+let Extension = concatExtension<BaseExt, "__opencl_c_atomic_order_seq_cst __opencl_c_atomic_scope_device">.ret in {
-PointerType<TypePair[1], GenericAS>, TypePair[1], MemoryOrder, MemoryOrder]>;
+def : Builtin<"atomic_compare_exchange_" # Variant,
-def : Builtin<"atomic_compare_exchange_" # Variant # "_explicit",
+[Bool, PointerType<VolatileType<TypePair[0]>, addrspace>,
-[Bool, PointerType<VolatileType<TypePair[0]>, GenericAS>,
+PointerType<TypePair[1], exp_ptr_addrspace>, TypePair[1]]>;
-PointerType<TypePair[1], GenericAS>, TypePair[1], MemoryOrder, MemoryOrder, MemoryScope]>;
+}
+let Extension = concatExtension<BaseExt, "__opencl_c_atomic_scope_device">.ret in {
+def : Builtin<"atomic_compare_exchange_" # Variant # "_explicit",
+[Bool, PointerType<VolatileType<TypePair[0]>, addrspace>,
+PointerType<TypePair[1], exp_ptr_addrspace>, TypePair[1], MemoryOrder, MemoryOrder]>;
+}
+let Extension = BaseExt in {
+def : Builtin<"atomic_compare_exchange_" # Variant # "_explicit",
+[Bool, PointerType<VolatileType<TypePair[0]>, addrspace>,
+PointerType<TypePair[1], exp_ptr_addrspace>, TypePair[1], MemoryOrder, MemoryOrder, MemoryScope]>;
+}
+}
 }
 }
 foreach TypePair = [[AtomicInt, Int, Int], [AtomicUInt, UInt, UInt],
 [AtomicLong, Long, Long], [AtomicULong, ULong, ULong],
 [AtomicUIntPtr, UIntPtr, PtrDiff]] in {
 foreach ModOp = ["add", "sub"] in {
-def : Builtin<"atomic_fetch_" # ModOp,
+defm : BuiltinAtomicExplicit<"atomic_fetch_" # ModOp,
-[TypePair[1], PointerType<VolatileType<TypePair[0]>, GenericAS>, TypePair[2]]>;
+[TypePair[1], PointerType<VolatileType<TypePair[0]>, addrspace>, TypePair[2]], BaseExt>;
-def : Builtin<"atomic_fetch_" # ModOp # "_explicit",
-[TypePair[1], PointerType<VolatileType<TypePair[0]>, GenericAS>, TypePair[2], MemoryOrder]>;
-def : Builtin<"atomic_fetch_" # ModOp # "_explicit",
-[TypePair[1], PointerType<VolatileType<TypePair[0]>, GenericAS>, TypePair[2], MemoryOrder, MemoryScope]>;
 }
 }
 foreach TypePair = [[AtomicInt, Int, Int], [AtomicUInt, UInt, UInt],
 [AtomicLong, Long, Long], [AtomicULong, ULong, ULong]] in {
 foreach ModOp = ["or", "xor", "and", "min", "max"] in {
-def : Builtin<"atomic_fetch_" # ModOp,
+defm : BuiltinAtomicExplicit<"atomic_fetch_" # ModOp,
-[TypePair[1], PointerType<VolatileType<TypePair[0]>, GenericAS>, TypePair[2]]>;
+[TypePair[1], PointerType<VolatileType<TypePair[0]>, addrspace>, TypePair[2]], BaseExt>;
-def : Builtin<"atomic_fetch_" # ModOp # "_explicit",
+}
-[TypePair[1], PointerType<VolatileType<TypePair[0]>, GenericAS>, TypePair[2], MemoryOrder]>;
+}
-def : Builtin<"atomic_fetch_" # ModOp # "_explicit",
-[TypePair[1], PointerType<VolatileType<TypePair[0]>, GenericAS>, TypePair[2], MemoryOrder, MemoryScope]>;
+defm : BuiltinAtomicExplicit<"atomic_flag_clear",
-}
+[Void, PointerType<VolatileType<AtomicFlag>, addrspace>], BaseExt>;
-}
+defm : BuiltinAtomicExplicit<"atomic_flag_test_and_set",
-def : Builtin<"atomic_flag_clear",
+[Bool, PointerType<VolatileType<AtomicFlag>, addrspace>], BaseExt>;
-[Void, PointerType<VolatileType<AtomicFlag>, GenericAS>]>;
+}
-def : Builtin<"atomic_flag_clear_explicit",
-[Void, PointerType<VolatileType<AtomicFlag>, GenericAS>, MemoryOrder]>;
+let MinVersion = CL20 in {
-def : Builtin<"atomic_flag_clear_explicit",
+def : Builtin<"atomic_work_item_fence", [Void, MemFenceFlags, MemoryOrder, MemoryScope]>;
-[Void, PointerType<VolatileType<AtomicFlag>, GenericAS>, MemoryOrder, MemoryScope]>;
+defm : OpenCL2Atomics<GenericAS, FuncExtOpenCLCGenericAddressSpace>;
-def : Builtin<"atomic_flag_test_and_set",
+defm : OpenCL2Atomics<GlobalAS, FuncExtOpenCLCNamedAddressSpaceBuiltins>;
-[Bool, PointerType<VolatileType<AtomicFlag>, GenericAS>]>;
+defm : OpenCL2Atomics<LocalAS, FuncExtOpenCLCNamedAddressSpaceBuiltins>;
-def : Builtin<"atomic_flag_test_and_set_explicit",
+}
-[Bool, PointerType<VolatileType<AtomicFlag>, GenericAS>, MemoryOrder]>;
-def : Builtin<"atomic_flag_test_and_set_explicit",
+// The functionality added by cl_ext_float_atomics extension
-[Bool, PointerType<VolatileType<AtomicFlag>, GenericAS>, MemoryOrder, MemoryScope]>;
+let MinVersion = CL20 in {
+foreach addrspace = [GlobalAS, LocalAS, GenericAS] in {
+defvar extension_fp16 = !cast<FunctionExtension>("FuncExtFloatAtomicsFp16" # addrspace # "LoadStore");
+defm : BuiltinAtomicExplicit<"atomic_store",
+[Void, PointerType<VolatileType<AtomicHalf>, addrspace>, AtomicHalf], extension_fp16>;
+defm : BuiltinAtomicExplicit<"atomic_load",
+[Half, PointerType<VolatileType<AtomicHalf>, addrspace>], extension_fp16>;
+defm : BuiltinAtomicExplicit<"atomic_exchange",
+[Half, PointerType<VolatileType<AtomicHalf>, addrspace>, Half], extension_fp16>;
+foreach ModOp = ["add", "sub"] in {
+defvar extension_fp16 = !cast<FunctionExtension>("FuncExtFloatAtomicsFp16" # addrspace # "Add");
+defvar extension_fp32 = !cast<FunctionExtension>("FuncExtFloatAtomicsFp32" # addrspace # "Add");
+defvar extension_fp64 = !cast<FunctionExtension>("FuncExtFloatAtomicsFp64" # addrspace # "Add");
+defm : BuiltinAtomicExplicit<"atomic_fetch_" # ModOp,
+[Half, PointerType<VolatileType<AtomicHalf>, addrspace>, Half], extension_fp16>;
+defm : BuiltinAtomicExplicit<"atomic_fetch_" # ModOp,
+[Float, PointerType<VolatileType<AtomicFloat>, addrspace>, Float], extension_fp32>;
+defm : BuiltinAtomicExplicit<"atomic_fetch_" # ModOp,
+[Double, PointerType<VolatileType<AtomicDouble>, addrspace>, Double], extension_fp64>;
+}
+foreach ModOp = ["min", "max"] in {
+defvar extension_fp16 = !cast<FunctionExtension>("FuncExtFloatAtomicsFp16" # addrspace # "MinMax");
+defvar extension_fp32 = !cast<FunctionExtension>("FuncExtFloatAtomicsFp32" # addrspace # "MinMax");
+defvar extension_fp64 = !cast<FunctionExtension>("FuncExtFloatAtomicsFp64" # addrspace # "MinMax");
+defm : BuiltinAtomicExplicit<"atomic_fetch_" # ModOp,
+[Half, PointerType<VolatileType<AtomicHalf>, addrspace>, Half], extension_fp16>;
+defm : BuiltinAtomicExplicit<"atomic_fetch_" # ModOp,
+[Float, PointerType<VolatileType<AtomicFloat>, addrspace>, Float], extension_fp32>;
+defm : BuiltinAtomicExplicit<"atomic_fetch_" # ModOp,
+[Double, PointerType<VolatileType<AtomicDouble>, addrspace>, Double], extension_fp64>;
+}
+}
 }
 //--------------------------------------------------------------------
 // OpenCL v1.1 s6.11.12, v1.2 s6.12.12, v2.0 s6.13.12 - Miscellaneous Vector Functions
 // --- Table 19 ---
 def : Builtin<"read_imagef", [Float, ImageType<Image2dDepth, "RO">, Sampler, VectorType<coordTy, 2>], Attr.Pure>;
 def : Builtin<"read_imagef", [Float, ImageType<Image2dArrayDepth, "RO">, Sampler, VectorType<coordTy, 4>], Attr.Pure>;
 }
 // --- Table 23: Sampler-less Read Functions ---
+multiclass ImageReadSamplerless<string aQual> {
+foreach imgTy = [Image2d, Image1dArray] in {
+def : Builtin<"read_imagef", [VectorType<Float, 4>, ImageType<imgTy, aQual>, VectorType<Int, 2>], Attr.Pure>;
+def : Builtin<"read_imagei", [VectorType<Int, 4>, ImageType<imgTy, aQual>, VectorType<Int, 2>], Attr.Pure>;
+def : Builtin<"read_imageui", [VectorType<UInt, 4>, ImageType<imgTy, aQual>, VectorType<Int, 2>], Attr.Pure>;
+}
+foreach imgTy = [Image3d, Image2dArray] in {
+def : Builtin<"read_imagef", [VectorType<Float, 4>, ImageType<imgTy, aQual>, VectorType<Int, 4>], Attr.Pure>;
+def : Builtin<"read_imagei", [VectorType<Int, 4>, ImageType<imgTy, aQual>, VectorType<Int, 4>], Attr.Pure>;
+def : Builtin<"read_imageui", [VectorType<UInt, 4>, ImageType<imgTy, aQual>, VectorType<Int, 4>], Attr.Pure>;
+}
+foreach imgTy = [Image1d, Image1dBuffer] in {
+def : Builtin<"read_imagef", [VectorType<Float, 4>, ImageType<imgTy, aQual>, Int], Attr.Pure>;
+def : Builtin<"read_imagei", [VectorType<Int, 4>, ImageType<imgTy, aQual>, Int], Attr.Pure>;
+def : Builtin<"read_imageui", [VectorType<UInt, 4>, ImageType<imgTy, aQual>, Int], Attr.Pure>;
+}
+def : Builtin<"read_imagef", [Float, ImageType<Image2dDepth, aQual>, VectorType<Int, 2>], Attr.Pure>;
+def : Builtin<"read_imagef", [Float, ImageType<Image2dArrayDepth, aQual>, VectorType<Int, 4>], Attr.Pure>;
+}
 let MinVersion = CL12 in {
-foreach aQual = ["RO", "RW"] in {
+defm : ImageReadSamplerless<"RO">;
-foreach imgTy = [Image2d, Image1dArray] in {
+let Extension = FuncExtOpenCLCReadWriteImages in {
-def : Builtin<"read_imagef", [VectorType<Float, 4>, ImageType<imgTy, aQual>, VectorType<Int, 2>], Attr.Pure>;
+defm : ImageReadSamplerless<"RW">;
-def : Builtin<"read_imagei", [VectorType<Int, 4>, ImageType<imgTy, aQual>, VectorType<Int, 2>], Attr.Pure>;
-def : Builtin<"read_imageui", [VectorType<UInt, 4>, ImageType<imgTy, aQual>, VectorType<Int, 2>], Attr.Pure>;
-}
-foreach imgTy = [Image3d, Image2dArray] in {
-def : Builtin<"read_imagef", [VectorType<Float, 4>, ImageType<imgTy, aQual>, VectorType<Int, 4>], Attr.Pure>;
-def : Builtin<"read_imagei", [VectorType<Int, 4>, ImageType<imgTy, aQual>, VectorType<Int, 4>], Attr.Pure>;
-def : Builtin<"read_imageui", [VectorType<UInt, 4>, ImageType<imgTy, aQual>, VectorType<Int, 4>], Attr.Pure>;
-}
-foreach imgTy = [Image1d, Image1dBuffer] in {
-def : Builtin<"read_imagef", [VectorType<Float, 4>, ImageType<imgTy, aQual>, Int], Attr.Pure>;
-def : Builtin<"read_imagei", [VectorType<Int, 4>, ImageType<imgTy, aQual>, Int], Attr.Pure>;
-def : Builtin<"read_imageui", [VectorType<UInt, 4>, ImageType<imgTy, aQual>, Int], Attr.Pure>;
-}
-def : Builtin<"read_imagef", [Float, ImageType<Image2dDepth, aQual>, VectorType<Int, 2>], Attr.Pure>;
-def : Builtin<"read_imagef", [Float, ImageType<Image2dArrayDepth, aQual>, VectorType<Int, 4>], Attr.Pure>;
 }
 }
 // --- Table 24: Image Write Functions ---
-foreach aQual = ["WO", "RW"] in {
+multiclass ImageWrite<string aQual> {
 foreach imgTy = [Image2d] in {
 def : Builtin<"write_imagef", [Void, ImageType<imgTy, aQual>, VectorType<Int, 2>, VectorType<Float, 4>]>;
 def : Builtin<"write_imagei", [Void, ImageType<imgTy, aQual>, VectorType<Int, 2>, VectorType<Int, 4>]>;
 def : Builtin<"write_imageui", [Void, ImageType<imgTy, aQual>, VectorType<Int, 2>, VectorType<UInt, 4>]>;
 }
 }
 def : Builtin<"write_imagef", [Void, ImageType<Image2dDepth, aQual>, VectorType<Int, 2>, Float]>;
 def : Builtin<"write_imagef", [Void, ImageType<Image2dArrayDepth, aQual>, VectorType<Int, 4>, Float]>;
 }
+defm : ImageWrite<"WO">;
+let Extension = FuncExtOpenCLCReadWriteImages in {
+defm : ImageWrite<"RW">;
+}
 // --- Table 25: Image Query Functions ---
-foreach aQual = ["RO", "WO", "RW"] in {
+multiclass ImageQuery<string aQual> {
 foreach imgTy = [Image1d, Image1dBuffer, Image2d, Image3d,
 Image1dArray, Image2dArray, Image2dDepth,
 Image2dArrayDepth] in {
 foreach name = ["get_image_width", "get_image_channel_data_type",
 "get_image_channel_order"] in {
 }
 def : Builtin<"get_image_dim", [VectorType<Int, 4>, ImageType<Image3d, aQual>], Attr.Const>;
 foreach imgTy = [Image1dArray, Image2dArray, Image2dArrayDepth] in {
 def : Builtin<"get_image_array_size", [Size, ImageType<imgTy, aQual>], Attr.Const>;
 }
+}
+defm : ImageQuery<"RO">;
+defm : ImageQuery<"WO">;
+let Extension = FuncExtOpenCLCReadWriteImages in {
+defm : ImageQuery<"RW">;
 }
 // OpenCL extension v2.0 s5.1.9: Built-in Image Read Functions
 // --- Table 8 ---
 foreach aQual = ["RO"] in {
 }
 }
 // OpenCL extension v2.0 s5.1.10: Built-in Image Sampler-less Read Functions
 // --- Table 9 ---
 let MinVersion = CL12 in {
-foreach aQual = ["RO", "RW"] in {
+multiclass ImageReadHalf<string aQual> {
 foreach name = ["read_imageh"] in {
 foreach imgTy = [Image2d, Image1dArray] in {
 def : Builtin<name, [VectorType<Half, 4>, ImageType<imgTy, aQual>, VectorType<Int, 2>], Attr.Pure>;
 }
 foreach imgTy = [Image3d, Image2dArray] in {
 foreach imgTy = [Image1d, Image1dBuffer] in {
 def : Builtin<name, [VectorType<Half, 4>, ImageType<imgTy, aQual>, Int], Attr.Pure>;
 }
 }
 }
+defm : ImageReadHalf<"RO">;
+let Extension = FuncExtOpenCLCReadWriteImages in {
+defm : ImageReadHalf<"RW">;
+}
 }
 // OpenCL extension v2.0 s5.1.11: Built-in Image Write Functions
 // --- Table 10 ---
-foreach aQual = ["WO", "RW"] in {
+multiclass ImageWriteHalf<string aQual> {
 foreach name = ["write_imageh"] in {
 def : Builtin<name, [Void, ImageType<Image2d, aQual>, VectorType<Int, 2>, VectorType<Half, 4>]>;
 def : Builtin<name, [Void, ImageType<Image2dArray, aQual>, VectorType<Int, 4>, VectorType<Half, 4>]>;
 def : Builtin<name, [Void, ImageType<Image1d, aQual>, Int, VectorType<Half, 4>]>;
 def : Builtin<name, [Void, ImageType<Image1dBuffer, aQual>, Int, VectorType<Half, 4>]>;
 def : Builtin<name, [Void, ImageType<Image1dArray, aQual>, VectorType<Int, 2>, VectorType<Half, 4>]>;
 def : Builtin<name, [Void, ImageType<Image3d, aQual>, VectorType<Int, 4>, VectorType<Half, 4>]>;
 }
 }
+defm : ImageWriteHalf<"WO">;
+let Extension = FuncExtOpenCLCReadWriteImages in {
+defm : ImageWriteHalf<"RW">;
+}
 //--------------------------------------------------------------------
 // OpenCL v2.0 s6.13.15 - Work-group Functions
 // --- Table 26 ---
-let MinVersion = CL20 in {
+let Extension = FuncExtOpenCLCWGCollectiveFunctions in {
 foreach name = ["work_group_all", "work_group_any"] in {
 def : Builtin<name, [Int, Int], Attr.Convergent>;
 }
 foreach name = ["work_group_broadcast"] in {
 def : Builtin<name, [IntLongFloatGenType1, IntLongFloatGenType1, Size], Attr.Convergent>;
 // --- Table 27 ---
 // Defined in Builtins.def
 // --- Table 28 ---
 // Builtins taking pipe arguments are defined in Builtins.def
-def : Builtin<"is_valid_reserve_id", [Bool, ReserveId]>;
+let Extension = FuncExtOpenCLCPipes in {
+def : Builtin<"is_valid_reserve_id", [Bool, ReserveId]>;
+}
 // --- Table 29 ---
 // Defined in Builtins.def
 // --- Table 32 ---
 // Defined in Builtins.def
 // --- Table 33 ---
-let MinVersion = CL20 in {
+let Extension = FuncExtOpenCLCDeviceEnqueue in {
 def : Builtin<"enqueue_marker",
 [Int, Queue, UInt, PointerType<ConstType<ClkEvent>, GenericAS>, PointerType<ClkEvent, GenericAS>]>;
 // --- Table 34 ---
 def : Builtin<"retain_event", [Void, ClkEvent]>;
 def : Builtin<name, [Float, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, Float], Attr.Pure>;
 def : Builtin<name, [Float, ImageType<imgTy, aQual>, Sampler, VectorType<Float, 4>, VectorType<Float, 2>, VectorType<Float, 2>], Attr.Pure>;
 }
 }
 }
-// Added to section 6.13.14.5
+}
-foreach aQual = ["RO", "WO", "RW"] in {
-foreach imgTy = [Image1d, Image2d, Image3d, Image1dArray, Image2dArray, Image2dDepth, Image2dArrayDepth] in {
+// Added to section 6.13.14.5
-def : Builtin<"get_image_num_mip_levels", [Int, ImageType<imgTy, aQual>]>;
+multiclass ImageQueryNumMipLevels<string aQual> {
-}
+foreach imgTy = [Image1d, Image2d, Image3d, Image1dArray, Image2dArray, Image2dDepth, Image2dArrayDepth] in {
-}
+def : Builtin<"get_image_num_mip_levels", [Int, ImageType<imgTy, aQual>]>;
+}
+}
+let Extension = FuncExtKhrMipmapImage in {
+defm : ImageQueryNumMipLevels<"RO">;
+defm : ImageQueryNumMipLevels<"WO">;
+defm : ImageQueryNumMipLevels<"RW">;
 }
 // Write functions are enabled using a separate extension.
 let Extension = FuncExtKhrMipmapImageWrites in {
 // Added to section 6.13.14.4.
 def : Builtin<"write_imagef", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int, VectorType<Float, 4>]>;
 def : Builtin<"write_imagei", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int, VectorType<Int, 4>]>;
 def : Builtin<"write_imageui", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int, VectorType<UInt, 4>]>;
 }
 def : Builtin<"write_imagef", [Void, ImageType<Image2dArrayDepth, aQual>, VectorType<Int, 4>, Int, Float]>;
-let Extension = FuncExtKhrMipmapWritesAndWrite3d in {
+foreach imgTy = [Image3d] in {
-foreach imgTy = [Image3d] in {
+def : Builtin<"write_imagef", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int, VectorType<Float, 4>]>;
-def : Builtin<"write_imagef", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int, VectorType<Float, 4>]>;
+def : Builtin<"write_imagei", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int, VectorType<Int, 4>]>;
-def : Builtin<"write_imagei", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int, VectorType<Int, 4>]>;
+def : Builtin<"write_imageui", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int, VectorType<UInt, 4>]>;
-def : Builtin<"write_imageui", [Void, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int, VectorType<UInt, 4>]>;
-}
 }
 }
 }
 //--------------------------------------------------------------------
 // OpenCL Extension v2.0 s18.3 - Creating OpenCL Memory Objects from OpenGL MSAA Textures
+// --- Table 6.13.14.3 ---
+multiclass ImageReadMsaa<string aQual> {
+foreach imgTy = [Image2dMsaa] in {
+def : Builtin<"read_imagef", [VectorType<Float, 4>, ImageType<imgTy, aQual>, VectorType<Int, 2>, Int], Attr.Pure>;
+def : Builtin<"read_imagei", [VectorType<Int, 4>, ImageType<imgTy, aQual>, VectorType<Int, 2>, Int], Attr.Pure>;
+def : Builtin<"read_imageui", [VectorType<UInt, 4>, ImageType<imgTy, aQual>, VectorType<Int, 2>, Int], Attr.Pure>;
+}
+foreach imgTy = [Image2dArrayMsaa] in {
+def : Builtin<"read_imagef", [VectorType<Float, 4>, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int], Attr.Pure>;
+def : Builtin<"read_imagei", [VectorType<Int, 4>, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int], Attr.Pure>;
+def : Builtin<"read_imageui", [VectorType<UInt, 4>, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int], Attr.Pure>;
+}
+foreach name = ["read_imagef"] in {
+def : Builtin<name, [Float, ImageType<Image2dMsaaDepth, aQual>, VectorType<Int, 2>, Int], Attr.Pure>;
+def : Builtin<name, [Float, ImageType<Image2dArrayMsaaDepth, aQual>, VectorType<Int, 4>, Int], Attr.Pure>;
+}
+}
+// --- Table 6.13.14.5 ---
+multiclass ImageQueryMsaa<string aQual> {
+foreach imgTy = [Image2dMsaa, Image2dArrayMsaa, Image2dMsaaDepth, Image2dArrayMsaaDepth] in {
+foreach name = ["get_image_width", "get_image_height",
+"get_image_channel_data_type", "get_image_channel_order",
+"get_image_num_samples"] in {
+def : Builtin<name, [Int, ImageType<imgTy, aQual>], Attr.Const>;
+}
+def : Builtin<"get_image_dim", [VectorType<Int, 2>, ImageType<imgTy, aQual>], Attr.Const>;
+}
+foreach imgTy = [Image2dArrayMsaa, Image2dArrayMsaaDepth] in {
+def : Builtin<"get_image_array_size", [Size, ImageType<imgTy, aQual>], Attr.Const>;
+}
+}
 let Extension = FuncExtKhrGlMsaaSharing in {
-// --- Table 6.13.14.3 ---
+defm : ImageReadMsaa<"RO">;
-foreach aQual = ["RO", "RW"] in {
+defm : ImageQueryMsaa<"RO">;
-foreach imgTy = [Image2dMsaa] in {
+defm : ImageQueryMsaa<"WO">;
-def : Builtin<"read_imagef", [VectorType<Float, 4>, ImageType<imgTy, aQual>, VectorType<Int, 2>, Int], Attr.Pure>;
+defm : ImageReadMsaa<"RW">;
-def : Builtin<"read_imagei", [VectorType<Int, 4>, ImageType<imgTy, aQual>, VectorType<Int, 2>, Int], Attr.Pure>;
+defm : ImageQueryMsaa<"RW">;
-def : Builtin<"read_imageui", [VectorType<UInt, 4>, ImageType<imgTy, aQual>, VectorType<Int, 2>, Int], Attr.Pure>;
-}
-foreach imgTy = [Image2dArrayMsaa] in {
-def : Builtin<"read_imagef", [VectorType<Float, 4>, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int], Attr.Pure>;
-def : Builtin<"read_imagei", [VectorType<Int, 4>, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int], Attr.Pure>;
-def : Builtin<"read_imageui", [VectorType<UInt, 4>, ImageType<imgTy, aQual>, VectorType<Int, 4>, Int], Attr.Pure>;
-}
-foreach name = ["read_imagef"] in {
-def : Builtin<name, [Float, ImageType<Image2dMsaaDepth, aQual>, VectorType<Int, 2>, Int], Attr.Pure>;
-def : Builtin<name, [Float, ImageType<Image2dArrayMsaaDepth, aQual>, VectorType<Int, 4>, Int], Attr.Pure>;
-}
-}
-// --- Table 6.13.14.5 ---
-foreach aQual = ["RO", "WO", "RW"] in {
-foreach imgTy = [Image2dMsaa, Image2dArrayMsaa, Image2dMsaaDepth, Image2dArrayMsaaDepth] in {
-foreach name = ["get_image_width", "get_image_height",
-"get_image_channel_data_type", "get_image_channel_order",
-"get_image_num_samples"] in {
-def : Builtin<name, [Int, ImageType<imgTy, aQual>], Attr.Const>;
-}
-def : Builtin<"get_image_dim", [VectorType<Int, 2>, ImageType<imgTy, aQual>], Attr.Const>;
-}
-foreach imgTy = [Image2dArrayMsaa, Image2dArrayMsaaDepth] in {
-def : Builtin<"get_image_array_size", [Size, ImageType<imgTy, aQual>], Attr.Const>;
-}
-}
 }
 //--------------------------------------------------------------------
 // OpenCL Extension v2.0 s28 - Subgroups
 // --- Table 28.2.1 ---
 }
 // --- Table 28.2.2 ---
 let Extension = FuncExtKhrSubgroups in {
 def : Builtin<"sub_group_barrier", [Void, MemFenceFlags], Attr.Convergent>;
-def : Builtin<"sub_group_barrier", [Void, MemFenceFlags, MemoryScope], Attr.Convergent>;
+let MinVersion = CL20 in {
+def : Builtin<"sub_group_barrier", [Void, MemFenceFlags, MemoryScope], Attr.Convergent>;
+}
 }
 // --- Table 28.2.4 ---
 let Extension = FuncExtKhrSubgroups in {
 foreach name = ["sub_group_all", "sub_group_any"] in {
 foreach op = ["and", "or", "xor"] in {
 def : Builtin<"sub_group_clustered_reduce_logical_" # op, [Int, Int, UInt]>;
 }
 }
+// Section 40.3.1 - cl_khr_extended_bit_ops
+let Extension = FuncExtKhrExtendedBitOps in {
+def : Builtin<"bitfield_insert", [AIGenTypeN, AIGenTypeN, AIGenTypeN, UInt, UInt], Attr.Const>;
+def : Builtin<"bitfield_extract_signed", [SGenTypeN, SGenTypeN, UInt, UInt], Attr.Const>;
+def : Builtin<"bitfield_extract_signed", [SGenTypeN, UGenTypeN, UInt, UInt], Attr.Const>;
+def : Builtin<"bitfield_extract_unsigned", [UGenTypeN, SGenTypeN, UInt, UInt], Attr.Const>;
+def : Builtin<"bitfield_extract_unsigned", [UGenTypeN, UGenTypeN, UInt, UInt], Attr.Const>;
+def : Builtin<"bit_reverse", [AIGenTypeN, AIGenTypeN], Attr.Const>;
+}
+// Section 42.3 - cl_khr_integer_dot_product
+let Extension = FunctionExtension<"__opencl_c_integer_dot_product_input_4x8bit"> in {
+def : Builtin<"dot", [UInt, VectorType<UChar, 4>, VectorType<UChar, 4>], Attr.Const>;
+def : Builtin<"dot", [Int, VectorType<Char, 4>, VectorType<Char, 4>], Attr.Const>;
+def : Builtin<"dot", [Int, VectorType<UChar, 4>, VectorType<Char, 4>], Attr.Const>;
+def : Builtin<"dot", [Int, VectorType<Char, 4>, VectorType<UChar, 4>], Attr.Const>;
+def : Builtin<"dot_acc_sat", [UInt, VectorType<UChar, 4>, VectorType<UChar, 4>, UInt], Attr.Const>;
+def : Builtin<"dot_acc_sat", [Int, VectorType<Char, 4>, VectorType<Char, 4>, Int], Attr.Const>;
+def : Builtin<"dot_acc_sat", [Int, VectorType<UChar, 4>, VectorType<Char, 4>, Int], Attr.Const>;
+def : Builtin<"dot_acc_sat", [Int, VectorType<Char, 4>, VectorType<UChar, 4>, Int], Attr.Const>;
+}
+let Extension = FunctionExtension<"__opencl_c_integer_dot_product_input_4x8bit_packed"> in {
+def : Builtin<"dot_4x8packed_uu_uint", [UInt, UInt, UInt], Attr.Const>;
+def : Builtin<"dot_4x8packed_ss_int", [Int, UInt, UInt], Attr.Const>;
+def : Builtin<"dot_4x8packed_us_int", [Int, UInt, UInt], Attr.Const>;
+def : Builtin<"dot_4x8packed_su_int", [Int, UInt, UInt], Attr.Const>;
+def : Builtin<"dot_acc_sat_4x8packed_uu_uint", [UInt, UInt, UInt, UInt], Attr.Const>;
+def : Builtin<"dot_acc_sat_4x8packed_ss_int", [Int, UInt, UInt, Int], Attr.Const>;
+def : Builtin<"dot_acc_sat_4x8packed_us_int", [Int, UInt, UInt, Int], Attr.Const>;
+def : Builtin<"dot_acc_sat_4x8packed_su_int", [Int, UInt, UInt, Int], Attr.Const>;
+}
+// Section 48.3 - cl_khr_subgroup_rotate
+let Extension = FunctionExtension<"cl_khr_subgroup_rotate"> in {
+def : Builtin<"sub_group_rotate", [AGenType1, AGenType1, Int], Attr.Convergent>;
+def : Builtin<"sub_group_clustered_rotate", [AGenType1, AGenType1, Int, UInt], Attr.Convergent>;
+}
 //--------------------------------------------------------------------
 // Arm extensions.
 let Extension = ArmIntegerDotProductInt8 in {
 foreach name = ["arm_dot"] in {
 def : Builtin<name, [UInt, VectorType<UChar, 4>, VectorType<UChar, 4>]>;

Mercurial > hg > CbC > CbC_llvm

comparison clang/lib/Sema/OpenCLBuiltins.td @ 236:c4bab56944e8 llvm-original