diff options
Diffstat (limited to 'clang/lib/Sema/OpenCLBuiltins.td')
| -rw-r--r-- | clang/lib/Sema/OpenCLBuiltins.td | 500 |
1 files changed, 188 insertions, 312 deletions
diff --git a/clang/lib/Sema/OpenCLBuiltins.td b/clang/lib/Sema/OpenCLBuiltins.td index cd704ba2df13..dc158454556a 100644 --- a/clang/lib/Sema/OpenCLBuiltins.td +++ b/clang/lib/Sema/OpenCLBuiltins.td @@ -57,14 +57,33 @@ class FunctionExtension<string _Ext> : AbstractExtension<_Ext>; // 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">; @@ -80,44 +99,40 @@ def FuncExtKhrLocalInt32ExtendedAtomics : FunctionExtension<"cl_khr_local_int32 def FuncExtKhrInt64BaseAtomics : FunctionExtension<"cl_khr_int64_base_atomics">; def FuncExtKhrInt64ExtendedAtomics : FunctionExtension<"cl_khr_int64_extended_atomics">; def FuncExtKhrMipmapImage : FunctionExtension<"cl_khr_mipmap_image">; -def FuncExtKhrMipmapImageReadWrite : FunctionExtension<"cl_khr_mipmap_image __opencl_c_read_write_images">; def FuncExtKhrMipmapImageWrites : FunctionExtension<"cl_khr_mipmap_image_writes">; def FuncExtKhrGlMsaaSharing : FunctionExtension<"cl_khr_gl_msaa_sharing">; -def FuncExtKhrGlMsaaSharingReadWrite : FunctionExtension<"cl_khr_gl_msaa_sharing __opencl_c_read_write_images">; +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 FuncExtFloatAtomicsFp16GlobalLoadStore : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_global_atomic_load_store">; -def FuncExtFloatAtomicsFp16LocalLoadStore : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_local_atomic_load_store">; -def FuncExtFloatAtomicsFp16GenericLoadStore : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_global_atomic_load_store __opencl_c_ext_fp16_local_atomic_load_store">; -def FuncExtFloatAtomicsFp16GlobalAdd : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_global_atomic_add">; -def FuncExtFloatAtomicsFp32GlobalAdd : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp32_global_atomic_add">; -def FuncExtFloatAtomicsFp64GlobalAdd : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp64_global_atomic_add">; -def FuncExtFloatAtomicsFp16LocalAdd : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_local_atomic_add">; -def FuncExtFloatAtomicsFp32LocalAdd : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp32_local_atomic_add">; -def FuncExtFloatAtomicsFp64LocalAdd : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp64_local_atomic_add">; -def FuncExtFloatAtomicsFp16GenericAdd : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_local_atomic_add __opencl_c_ext_fp16_global_atomic_add">; -def FuncExtFloatAtomicsFp32GenericAdd : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp32_local_atomic_add __opencl_c_ext_fp32_global_atomic_add">; -def FuncExtFloatAtomicsFp64GenericAdd : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp64_local_atomic_add __opencl_c_ext_fp64_global_atomic_add">; -def FuncExtFloatAtomicsFp16GlobalMinMax : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_global_atomic_min_max">; -def FuncExtFloatAtomicsFp32GlobalMinMax : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp32_global_atomic_min_max">; -def FuncExtFloatAtomicsFp64GlobalMinMax : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp64_global_atomic_min_max">; -def FuncExtFloatAtomicsFp16LocalMinMax : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_local_atomic_min_max">; -def FuncExtFloatAtomicsFp32LocalMinMax : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp32_local_atomic_min_max">; -def FuncExtFloatAtomicsFp64LocalMinMax : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp64_local_atomic_min_max">; -def FuncExtFloatAtomicsFp16GenericMinMax : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp16_local_atomic_min_max __opencl_c_ext_fp16_global_atomic_min_max">; -def FuncExtFloatAtomicsFp32GenericMinMax : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp32_local_atomic_min_max __opencl_c_ext_fp32_global_atomic_min_max">; -def FuncExtFloatAtomicsFp64GenericMinMax : FunctionExtension<"cl_ext_float_atomics __opencl_c_ext_fp64_local_atomic_min_max __opencl_c_ext_fp64_global_atomic_min_max">; +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">; @@ -227,7 +242,12 @@ class ImageType<Type _Ty, string _AccessQualifier> : 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">, + !eq(_AccessQualifier, "RW") : TypeExtension<"__opencl_c_read_write_images">, + true : _Ty.Extension); } // OpenCL enum type (e.g. memory_scope). @@ -332,9 +352,22 @@ 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()">>; @@ -372,10 +405,14 @@ 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)">>; -def AtomicULong : Type<"atomic_ulong", QualType<"Context.getAtomicType(Context.UnsignedLongTy)">>; +let Extension = Atomic64TypeExt in { + 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())">>; @@ -853,22 +890,22 @@ defm : VloadVstore<[ConstantAS], 0>; multiclass VloadVstoreHalf<list<AddressSpace> addrspaces, bit defStores> { foreach AS = addrspaces in { - def : Builtin<"vload_half", [Float, Size, PointerType<ConstType<Half>, AS>], Attr.Pure>; + 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, PointerType<ConstType<Half>, AS>], Attr.Pure>; + 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 = ["vstore_half" # rnd] in { - def : Builtin<name, [Void, Float, Size, PointerType<Half, AS>]>; - def : Builtin<name, [Void, Double, Size, PointerType<Half, AS>]>; + def : Builtin<name, [Void, Float, Size, !cast<Type>("HalfPtr" # 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, PointerType<Half, AS>]>; - def : Builtin<name, [Void, VectorType<Double, VSize>, Size, PointerType<Half, AS>]>; + def : Builtin<name, [Void, VectorType<Float, VSize>, Size, !cast<Type>("HalfPtr" # AS)]>; + def : Builtin<name, [Void, VectorType<Double, VSize>, Size, !cast<Type>("HalfPtr" # AS)]>; } } } @@ -900,7 +937,7 @@ 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. @@ -1041,42 +1078,61 @@ let Extension = FuncExtOpenCLCxx in { } // 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", - [Void, PointerType<VolatileType<TypePair[0]>, GenericAS>, TypePair[1]]>; - def : Builtin<"atomic_store", - [Void, PointerType<VolatileType<TypePair[0]>, GenericAS>, TypePair[1]]>; - def : Builtin<"atomic_store_explicit", - [Void, PointerType<VolatileType<TypePair[0]>, GenericAS>, TypePair[1], MemoryOrder]>; - def : Builtin<"atomic_store_explicit", - [Void, PointerType<VolatileType<TypePair[0]>, GenericAS>, TypePair[1], MemoryOrder, MemoryScope]>; - def : Builtin<"atomic_load", - [TypePair[1], PointerType<VolatileType<TypePair[0]>, GenericAS>]>; - 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]>; + let Extension = BaseExt in { + def : Builtin<"atomic_init", + [Void, PointerType<VolatileType<TypePair[0]>, addrspace>, TypePair[1]]>; + } + defm : BuiltinAtomicExplicit<"atomic_store", + [Void, PointerType<VolatileType<TypePair[0]>, addrspace>, TypePair[1]], BaseExt>; + defm : BuiltinAtomicExplicit<"atomic_load", + [TypePair[1], PointerType<VolatileType<TypePair[0]>, addrspace>], BaseExt>; + defm : BuiltinAtomicExplicit<"atomic_exchange", + [TypePair[1], PointerType<VolatileType<TypePair[0]>, addrspace>, TypePair[1]], BaseExt>; foreach Variant = ["weak", "strong"] in { - def : Builtin<"atomic_compare_exchange_" # Variant, - [Bool, PointerType<VolatileType<TypePair[0]>, GenericAS>, - PointerType<TypePair[1], GenericAS>, TypePair[1]]>; - def : Builtin<"atomic_compare_exchange_" # Variant # "_explicit", - [Bool, PointerType<VolatileType<TypePair[0]>, GenericAS>, - PointerType<TypePair[1], GenericAS>, TypePair[1], MemoryOrder, MemoryOrder]>; - def : Builtin<"atomic_compare_exchange_" # Variant # "_explicit", - [Bool, PointerType<VolatileType<TypePair[0]>, GenericAS>, - PointerType<TypePair[1], GenericAS>, TypePair[1], MemoryOrder, MemoryOrder, MemoryScope]>; + foreach exp_ptr_addrspace = !cond( + !eq(BaseExt, FuncExtOpenCLCGenericAddressSpace): [GenericAS], + !eq(BaseExt, FuncExtOpenCLCNamedAddressSpaceBuiltins): [GlobalAS, LocalAS, PrivateAS]) + in { + let Extension = concatExtension<BaseExt, "__opencl_c_atomic_order_seq_cst __opencl_c_atomic_scope_device">.ret in { + def : Builtin<"atomic_compare_exchange_" # Variant, + [Bool, PointerType<VolatileType<TypePair[0]>, addrspace>, + PointerType<TypePair[1], exp_ptr_addrspace>, TypePair[1]]>; + } + 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]>; + } + } } } @@ -1084,249 +1140,69 @@ let MinVersion = CL20 in { [AtomicLong, Long, Long], [AtomicULong, ULong, ULong], [AtomicUIntPtr, UIntPtr, PtrDiff]] in { foreach ModOp = ["add", "sub"] in { - def : Builtin<"atomic_fetch_" # ModOp, - [TypePair[1], PointerType<VolatileType<TypePair[0]>, GenericAS>, TypePair[2]]>; - 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_fetch_" # ModOp, + [TypePair[1], PointerType<VolatileType<TypePair[0]>, addrspace>, TypePair[2]], BaseExt>; } } 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, - [TypePair[1], PointerType<VolatileType<TypePair[0]>, GenericAS>, TypePair[2]]>; - 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_fetch_" # ModOp, + [TypePair[1], PointerType<VolatileType<TypePair[0]>, addrspace>, TypePair[2]], BaseExt>; } } - def : Builtin<"atomic_flag_clear", - [Void, PointerType<VolatileType<AtomicFlag>, GenericAS>]>; - def : Builtin<"atomic_flag_clear_explicit", - [Void, PointerType<VolatileType<AtomicFlag>, GenericAS>, MemoryOrder]>; - def : Builtin<"atomic_flag_clear_explicit", - [Void, PointerType<VolatileType<AtomicFlag>, GenericAS>, MemoryOrder, MemoryScope]>; + defm : BuiltinAtomicExplicit<"atomic_flag_clear", + [Void, PointerType<VolatileType<AtomicFlag>, addrspace>], BaseExt>; - def : Builtin<"atomic_flag_test_and_set", - [Bool, PointerType<VolatileType<AtomicFlag>, GenericAS>]>; - def : Builtin<"atomic_flag_test_and_set_explicit", - [Bool, PointerType<VolatileType<AtomicFlag>, GenericAS>, MemoryOrder]>; - def : Builtin<"atomic_flag_test_and_set_explicit", - [Bool, PointerType<VolatileType<AtomicFlag>, GenericAS>, MemoryOrder, MemoryScope]>; + defm : BuiltinAtomicExplicit<"atomic_flag_test_and_set", + [Bool, PointerType<VolatileType<AtomicFlag>, addrspace>], BaseExt>; +} + +let MinVersion = CL20 in { + def : Builtin<"atomic_work_item_fence", [Void, MemFenceFlags, MemoryOrder, MemoryScope]>; + + defm : OpenCL2Atomics<GenericAS, FuncExtOpenCLCGenericAddressSpace>; + defm : OpenCL2Atomics<GlobalAS, FuncExtOpenCLCNamedAddressSpaceBuiltins>; + defm : OpenCL2Atomics<LocalAS, FuncExtOpenCLCNamedAddressSpaceBuiltins>; } // The functionality added by cl_ext_float_atomics extension let MinVersion = CL20 in { - let Extension = FuncExtFloatAtomicsFp16GlobalLoadStore in { - def : Builtin<"atomic_store", - [Void, PointerType<VolatileType<AtomicHalf>, GlobalAS>, AtomicHalf]>; - def : Builtin<"atomic_store_explicit", - [Void, PointerType<VolatileType<AtomicHalf>, GlobalAS>, AtomicHalf, MemoryOrder]>; - def : Builtin<"atomic_store_explicit", - [Void, PointerType<VolatileType<AtomicHalf>, GlobalAS>, AtomicHalf, MemoryOrder, MemoryScope]>; - def : Builtin<"atomic_load", - [Half, PointerType<VolatileType<AtomicHalf>, GlobalAS>]>; - def : Builtin<"atomic_load_explicit", - [Half, PointerType<VolatileType<AtomicHalf>, GlobalAS>, MemoryOrder]>; - def : Builtin<"atomic_load_explicit", - [Half, PointerType<VolatileType<AtomicHalf>, GlobalAS>, MemoryOrder, MemoryScope]>; - def : Builtin<"atomic_exchange", - [Half, PointerType<VolatileType<AtomicHalf>, GlobalAS>, Half]>; - def : Builtin<"atomic_exchange_explicit", - [Half, PointerType<VolatileType<AtomicHalf>, GlobalAS>, Half, MemoryOrder]>; - def : Builtin<"atomic_exchange_explicit", - [Half, PointerType<VolatileType<AtomicHalf>, GlobalAS>, Half, MemoryOrder, MemoryScope]>; - } - let Extension = FuncExtFloatAtomicsFp16LocalLoadStore in { - def : Builtin<"atomic_store", - [Void, PointerType<VolatileType<AtomicHalf>, LocalAS>, AtomicHalf]>; - def : Builtin<"atomic_store_explicit", - [Void, PointerType<VolatileType<AtomicHalf>, LocalAS>, AtomicHalf, MemoryOrder]>; - def : Builtin<"atomic_store_explicit", - [Void, PointerType<VolatileType<AtomicHalf>, LocalAS>, AtomicHalf, MemoryOrder, MemoryScope]>; - def : Builtin<"atomic_load", - [Half, PointerType<VolatileType<AtomicHalf>, LocalAS>]>; - def : Builtin<"atomic_load_explicit", - [Half, PointerType<VolatileType<AtomicHalf>, LocalAS>, MemoryOrder]>; - def : Builtin<"atomic_load_explicit", - [Half, PointerType<VolatileType<AtomicHalf>, LocalAS>, MemoryOrder, MemoryScope]>; - def : Builtin<"atomic_exchange", - [Half, PointerType<VolatileType<AtomicHalf>, LocalAS>, Half]>; - def : Builtin<"atomic_exchange_explicit", - [Half, PointerType<VolatileType<AtomicHalf>, LocalAS>, Half, MemoryOrder]>; - def : Builtin<"atomic_exchange_explicit", - [Half, PointerType<VolatileType<AtomicHalf>, LocalAS>, Half, MemoryOrder, MemoryScope]>; - } - let Extension = FuncExtFloatAtomicsFp16GenericLoadStore in { - def : Builtin<"atomic_store", - [Void, PointerType<VolatileType<AtomicHalf>, GenericAS>, AtomicHalf]>; - def : Builtin<"atomic_store_explicit", - [Void, PointerType<VolatileType<AtomicHalf>, GenericAS>, AtomicHalf, MemoryOrder]>; - def : Builtin<"atomic_store_explicit", - [Void, PointerType<VolatileType<AtomicHalf>, GenericAS>, AtomicHalf, MemoryOrder, MemoryScope]>; - def : Builtin<"atomic_load", - [Half, PointerType<VolatileType<AtomicHalf>, GenericAS>]>; - def : Builtin<"atomic_load_explicit", - [Half, PointerType<VolatileType<AtomicHalf>, GenericAS>, MemoryOrder]>; - def : Builtin<"atomic_load_explicit", - [Half, PointerType<VolatileType<AtomicHalf>, GenericAS>, MemoryOrder, MemoryScope]>; - def : Builtin<"atomic_exchange", - [Half, PointerType<VolatileType<AtomicHalf>, GenericAS>, Half]>; - def : Builtin<"atomic_exchange_explicit", - [Half, PointerType<VolatileType<AtomicHalf>, GenericAS>, Half, MemoryOrder]>; - def : Builtin<"atomic_exchange_explicit", - [Half, PointerType<VolatileType<AtomicHalf>, GenericAS>, Half, MemoryOrder, MemoryScope]>; - } - foreach ModOp = ["add", "sub"] in { - let Extension = FuncExtFloatAtomicsFp16GlobalAdd in { - def : Builtin<"atomic_fetch_" # ModOp, - [Half, PointerType<VolatileType<AtomicFloat>, GlobalAS>, Half]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Half, PointerType<VolatileType<AtomicFloat>, GlobalAS>, Half, MemoryOrder]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Half, PointerType<VolatileType<AtomicFloat>, GlobalAS>, Half, MemoryOrder, MemoryScope]>; - } - let Extension = FuncExtFloatAtomicsFp32GlobalAdd in { - def : Builtin<"atomic_fetch_" # ModOp, - [Float, PointerType<VolatileType<AtomicFloat>, GlobalAS>, Float]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Float, PointerType<VolatileType<AtomicFloat>, GlobalAS>, Float, MemoryOrder]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Float, PointerType<VolatileType<AtomicFloat>, GlobalAS>, Float, MemoryOrder, MemoryScope]>; - } - let Extension = FuncExtFloatAtomicsFp64GlobalAdd in { - def : Builtin<"atomic_fetch_" # ModOp, - [Double, PointerType<VolatileType<AtomicDouble>, GlobalAS>, Double]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Double, PointerType<VolatileType<AtomicDouble>, GlobalAS>, Double, MemoryOrder]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Double, PointerType<VolatileType<AtomicDouble>, GlobalAS>, Double, MemoryOrder, MemoryScope]>; - } - let Extension = FuncExtFloatAtomicsFp16LocalAdd in { - def : Builtin<"atomic_fetch_" # ModOp, - [Half, PointerType<VolatileType<AtomicFloat>, LocalAS>, Half]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Half, PointerType<VolatileType<AtomicFloat>, LocalAS>, Half, MemoryOrder]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Half, PointerType<VolatileType<AtomicFloat>, LocalAS>, Half, MemoryOrder, MemoryScope]>; - } - let Extension = FuncExtFloatAtomicsFp32LocalAdd in { - def : Builtin<"atomic_fetch_" # ModOp, - [Float, PointerType<VolatileType<AtomicFloat>, LocalAS>, Float]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Float, PointerType<VolatileType<AtomicFloat>, LocalAS>, Float, MemoryOrder]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Float, PointerType<VolatileType<AtomicFloat>, LocalAS>, Float, MemoryOrder, MemoryScope]>; - } - let Extension = FuncExtFloatAtomicsFp64LocalAdd in { - def : Builtin<"atomic_fetch_" # ModOp, - [Double, PointerType<VolatileType<AtomicDouble>, LocalAS>, Double]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Double, PointerType<VolatileType<AtomicDouble>, LocalAS>, Double, MemoryOrder]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Double, PointerType<VolatileType<AtomicDouble>, LocalAS>, Double, MemoryOrder, MemoryScope]>; - } - let Extension = FuncExtFloatAtomicsFp16GenericAdd in { - def : Builtin<"atomic_fetch_" # ModOp, - [Half, PointerType<VolatileType<AtomicFloat>, GenericAS>, Half]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Half, PointerType<VolatileType<AtomicFloat>, GenericAS>, Half, MemoryOrder]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Half, PointerType<VolatileType<AtomicFloat>, GenericAS>, Half, MemoryOrder, MemoryScope]>; - } - let Extension = FuncExtFloatAtomicsFp32GenericAdd in { - def : Builtin<"atomic_fetch_" # ModOp, - [Float, PointerType<VolatileType<AtomicFloat>, GenericAS>, Float]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Float, PointerType<VolatileType<AtomicFloat>, GenericAS>, Float, MemoryOrder]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Float, PointerType<VolatileType<AtomicFloat>, GenericAS>, Float, MemoryOrder, MemoryScope]>; - } - let Extension = FuncExtFloatAtomicsFp64GenericAdd in { - def : Builtin<"atomic_fetch_" # ModOp, - [Double, PointerType<VolatileType<AtomicDouble>, GenericAS>, Double]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Double, PointerType<VolatileType<AtomicDouble>, GenericAS>, Double, MemoryOrder]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Double, PointerType<VolatileType<AtomicDouble>, GenericAS>, Double, MemoryOrder, MemoryScope]>; - } - } - foreach ModOp = ["min", "max"] in { - let Extension = FuncExtFloatAtomicsFp16GlobalMinMax in { - def : Builtin<"atomic_fetch_" # ModOp, - [Half, PointerType<VolatileType<AtomicHalf>, GlobalAS>, Half]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Half, PointerType<VolatileType<AtomicHalf>, GlobalAS>, Half, MemoryOrder]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Half, PointerType<VolatileType<AtomicHalf>, GlobalAS>, Half, MemoryOrder, MemoryScope]>; - } - let Extension = FuncExtFloatAtomicsFp32GlobalMinMax in { - def : Builtin<"atomic_fetch_" # ModOp, - [Float, PointerType<VolatileType<AtomicFloat>, GlobalAS>, Float]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Float, PointerType<VolatileType<AtomicFloat>, GlobalAS>, Float, MemoryOrder]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Float, PointerType<VolatileType<AtomicFloat>, GlobalAS>, Float, MemoryOrder, MemoryScope]>; - } - let Extension = FuncExtFloatAtomicsFp64GlobalMinMax in { - def : Builtin<"atomic_fetch_" # ModOp, - [Double, PointerType<VolatileType<AtomicDouble>, GlobalAS>, Double]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Double, PointerType<VolatileType<AtomicDouble>, GlobalAS>, Double, MemoryOrder]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Double, PointerType<VolatileType<AtomicDouble>, GlobalAS>, Double, MemoryOrder, MemoryScope]>; - } - let Extension = FuncExtFloatAtomicsFp16LocalMinMax in { - def : Builtin<"atomic_fetch_" # ModOp, - [Half, PointerType<VolatileType<AtomicHalf>, LocalAS>, Half]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Half, PointerType<VolatileType<AtomicHalf>, LocalAS>, Half, MemoryOrder]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Half, PointerType<VolatileType<AtomicHalf>, LocalAS>, Half, MemoryOrder, MemoryScope]>; - } - let Extension = FuncExtFloatAtomicsFp32LocalMinMax in { - def : Builtin<"atomic_fetch_" # ModOp, - [Float, PointerType<VolatileType<AtomicFloat>, LocalAS>, Float]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Float, PointerType<VolatileType<AtomicFloat>, LocalAS>, Float, MemoryOrder]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Float, PointerType<VolatileType<AtomicFloat>, LocalAS>, Float, MemoryOrder, MemoryScope]>; - } - let Extension = FuncExtFloatAtomicsFp64LocalMinMax in { - def : Builtin<"atomic_fetch_" # ModOp, - [Double, PointerType<VolatileType<AtomicDouble>, LocalAS>, Double]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Double, PointerType<VolatileType<AtomicDouble>, LocalAS>, Double, MemoryOrder]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Double, PointerType<VolatileType<AtomicDouble>, LocalAS>, Double, MemoryOrder, MemoryScope]>; - } - let Extension = FuncExtFloatAtomicsFp16GenericMinMax in { - def : Builtin<"atomic_fetch_" # ModOp, - [Half, PointerType<VolatileType<AtomicHalf>, GenericAS>, Half]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Half, PointerType<VolatileType<AtomicHalf>, GenericAS>, Half, MemoryOrder]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Half, PointerType<VolatileType<AtomicHalf>, GenericAS>, Half, MemoryOrder, MemoryScope]>; - } - let Extension = FuncExtFloatAtomicsFp32GenericMinMax in { - def : Builtin<"atomic_fetch_" # ModOp, - [Float, PointerType<VolatileType<AtomicFloat>, GenericAS>, Float]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Float, PointerType<VolatileType<AtomicFloat>, GenericAS>, Float, MemoryOrder]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Float, PointerType<VolatileType<AtomicFloat>, GenericAS>, Float, MemoryOrder, MemoryScope]>; + 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>; } - let Extension = FuncExtFloatAtomicsFp64GenericMinMax in { - def : Builtin<"atomic_fetch_" # ModOp, - [Double, PointerType<VolatileType<AtomicDouble>, GenericAS>, Double]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Double, PointerType<VolatileType<AtomicDouble>, GenericAS>, Double, MemoryOrder]>; - def : Builtin<"atomic_fetch_" # ModOp # "_explicit", - [Double, PointerType<VolatileType<AtomicDouble>, GenericAS>, Double, MemoryOrder, MemoryScope]>; + + 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>; } } } @@ -1592,7 +1468,7 @@ let Extension = FuncExtOpenCLCPipes in { // 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>]>; @@ -1731,9 +1607,6 @@ multiclass ImageQueryNumMipLevels<string aQual> { let Extension = FuncExtKhrMipmapImage in { defm : ImageQueryNumMipLevels<"RO">; defm : ImageQueryNumMipLevels<"WO">; -} - -let Extension = FuncExtKhrMipmapImageReadWrite in { defm : ImageQueryNumMipLevels<"RW">; } @@ -1763,12 +1636,10 @@ let Extension = FuncExtKhrMipmapImageWrites in { 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 { - 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>]>; - } + foreach imgTy = [Image3d] in { + 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>]>; } } } @@ -1812,9 +1683,6 @@ let Extension = FuncExtKhrGlMsaaSharing in { defm : ImageReadMsaa<"RO">; defm : ImageQueryMsaa<"RO">; defm : ImageQueryMsaa<"WO">; -} - -let Extension = FuncExtKhrGlMsaaSharingReadWrite in { defm : ImageReadMsaa<"RW">; defm : ImageQueryMsaa<"RW">; } @@ -1838,7 +1706,9 @@ let Extension = FuncExtKhrSubgroups in { // --- 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 --- @@ -1975,6 +1845,12 @@ let Extension = FunctionExtension<"__opencl_c_integer_dot_product_input_4x8bit_p 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 { |