diff options
Diffstat (limited to 'contrib/llvm-project/clang/lib/Sema/OpenCLBuiltins.td')
| -rw-r--r-- | contrib/llvm-project/clang/lib/Sema/OpenCLBuiltins.td | 650 |
1 files changed, 364 insertions, 286 deletions
diff --git a/contrib/llvm-project/clang/lib/Sema/OpenCLBuiltins.td b/contrib/llvm-project/clang/lib/Sema/OpenCLBuiltins.td index cd704fe395a9..0cceba090bd8 100644 --- a/contrib/llvm-project/clang/lib/Sema/OpenCLBuiltins.td +++ b/contrib/llvm-project/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">; @@ -83,12 +102,37 @@ def FuncExtKhrMipmapImage : FunctionExtension<"cl_khr_mipmap_imag 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">; @@ -198,7 +242,13 @@ 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">, + !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). @@ -303,9 +353,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()">>; @@ -343,10 +406,15 @@ 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())">>; def AtomicSize : Type<"atomic_size_t", QualType<"Context.getAtomicType(Context.getSizeType())">>; @@ -543,9 +611,10 @@ foreach name = ["fma", "mad"] in { def : Builtin<name, [FGenTypeN, FGenTypeN, FGenTypeN, FGenTypeN], Attr.Const>; } -// --- Version dependent --- -let MaxVersion = CL20 in { - foreach AS = [GlobalAS, LocalAS, PrivateAS] in { +// The following math builtins take pointer arguments. Which overloads are +// available depends on whether the generic address space feature is enabled. +multiclass MathWithPointer<list<AddressSpace> addrspaces> { + foreach AS = addrspaces in { foreach name = ["fract", "modf", "sincos"] in { def : Builtin<name, [FGenTypeN, FGenTypeN, PointerType<FGenTypeN, AS>]>; } @@ -561,19 +630,12 @@ let MaxVersion = CL20 in { } } } -let MinVersion = CL20 in { - foreach name = ["fract", "modf", "sincos"] in { - def : Builtin<name, [FGenTypeN, FGenTypeN, PointerType<FGenTypeN, GenericAS>]>; - } - foreach name = ["frexp", "lgamma_r"] in { - foreach Type = [GenTypeFloatVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeHalfVecAndScalar] in { - 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>]>; - } - } + +let Extension = FuncExtOpenCLCNamedAddressSpaceBuiltins in { + defm : MathWithPointer<[GlobalAS, LocalAS, PrivateAS]>; +} +let Extension = FuncExtOpenCLCGenericAddressSpace in { + defm : MathWithPointer<[GenericAS]>; } // --- Table 9 --- @@ -783,165 +845,83 @@ foreach name = ["select"] in { // 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 -// and private address spaces. -let MaxVersion = CL20 in { - foreach AS = [GlobalAS, LocalAS, PrivateAS] in { +multiclass VloadVstore<list<AddressSpace> addrspaces, bit defStores> { + foreach AS = addrspaces 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<UChar, VSize>, Size, PointerType<ConstType<UChar>, AS>]>; - def : Builtin<name, [VectorType<Short, VSize>, Size, PointerType<ConstType<Short>, AS>]>; - def : Builtin<name, [VectorType<UShort, VSize>, Size, PointerType<ConstType<UShort>, AS>]>; - def : Builtin<name, [VectorType<Int, VSize>, Size, PointerType<ConstType<Int>, AS>]>; - def : Builtin<name, [VectorType<UInt, VSize>, Size, PointerType<ConstType<UInt>, AS>]>; - def : Builtin<name, [VectorType<Long, VSize>, Size, PointerType<ConstType<Long>, AS>]>; - def : Builtin<name, [VectorType<ULong, VSize>, Size, PointerType<ConstType<ULong>, AS>]>; - def : Builtin<name, [VectorType<Float, VSize>, Size, PointerType<ConstType<Float>, AS>]>; - def : Builtin<name, [VectorType<Double, VSize>, Size, PointerType<ConstType<Double>, AS>]>; - def : Builtin<name, [VectorType<Half, VSize>, Size, PointerType<ConstType<Half>, AS>]>; - } - foreach name = ["vstore" # VSize] in { - def : Builtin<name, [Void, VectorType<Char, VSize>, Size, PointerType<Char, AS>]>; - def : Builtin<name, [Void, VectorType<UChar, VSize>, Size, PointerType<UChar, AS>]>; - def : Builtin<name, [Void, VectorType<Short, VSize>, Size, PointerType<Short, AS>]>; - def : Builtin<name, [Void, VectorType<UShort, VSize>, Size, PointerType<UShort, AS>]>; - def : Builtin<name, [Void, VectorType<Int, VSize>, Size, PointerType<Int, AS>]>; - def : Builtin<name, [Void, VectorType<UInt, VSize>, Size, PointerType<UInt, AS>]>; - def : Builtin<name, [Void, VectorType<Long, VSize>, Size, PointerType<Long, AS>]>; - def : Builtin<name, [Void, VectorType<ULong, VSize>, Size, PointerType<ULong, AS>]>; - def : Builtin<name, [Void, VectorType<Float, VSize>, Size, PointerType<Float, 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>]>; + def : Builtin<name, [VectorType<Char, VSize>, Size, PointerType<ConstType<Char>, AS>], Attr.Pure>; + def : Builtin<name, [VectorType<UChar, VSize>, Size, PointerType<ConstType<UChar>, AS>], Attr.Pure>; + def : Builtin<name, [VectorType<Short, VSize>, Size, PointerType<ConstType<Short>, AS>], Attr.Pure>; + def : Builtin<name, [VectorType<UShort, VSize>, Size, PointerType<ConstType<UShort>, AS>], Attr.Pure>; + def : Builtin<name, [VectorType<Int, VSize>, Size, PointerType<ConstType<Int>, AS>], Attr.Pure>; + def : Builtin<name, [VectorType<UInt, VSize>, Size, PointerType<ConstType<UInt>, AS>], Attr.Pure>; + def : Builtin<name, [VectorType<Long, VSize>, Size, PointerType<ConstType<Long>, AS>], Attr.Pure>; + def : Builtin<name, [VectorType<ULong, VSize>, Size, PointerType<ConstType<ULong>, AS>], Attr.Pure>; + def : Builtin<name, [VectorType<Float, VSize>, Size, PointerType<ConstType<Float>, AS>], Attr.Pure>; + def : Builtin<name, [VectorType<Double, VSize>, Size, PointerType<ConstType<Double>, AS>], Attr.Pure>; + def : Builtin<name, [VectorType<Half, VSize>, Size, PointerType<ConstType<Half>, AS>], Attr.Pure>; } - foreach rnd = ["", "_rte", "_rtz", "_rtp", "_rtn"] in { - foreach name = ["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>]>; + if defStores then { + foreach name = ["vstore" # VSize] in { + def : Builtin<name, [Void, VectorType<Char, VSize>, Size, PointerType<Char, AS>]>; + def : Builtin<name, [Void, VectorType<UChar, VSize>, Size, PointerType<UChar, AS>]>; + def : Builtin<name, [Void, VectorType<Short, VSize>, Size, PointerType<Short, AS>]>; + def : Builtin<name, [Void, VectorType<UShort, VSize>, Size, PointerType<UShort, AS>]>; + def : Builtin<name, [Void, VectorType<Int, VSize>, Size, PointerType<Int, AS>]>; + def : Builtin<name, [Void, VectorType<UInt, VSize>, Size, PointerType<UInt, AS>]>; + def : Builtin<name, [Void, VectorType<Long, VSize>, Size, PointerType<Long, AS>]>; + def : Builtin<name, [Void, VectorType<ULong, VSize>, Size, PointerType<ULong, AS>]>; + def : Builtin<name, [Void, VectorType<Float, VSize>, Size, PointerType<Float, AS>]>; + def : Builtin<name, [Void, VectorType<Double, VSize>, Size, PointerType<Double, AS>]>; + def : Builtin<name, [Void, VectorType<Half, VSize>, Size, PointerType<Half, 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>]>; - } - } - } + +let Extension = FuncExtOpenCLCNamedAddressSpaceBuiltins in { + defm : VloadVstore<[GlobalAS, LocalAS, PrivateAS], 1>; } -// 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 Extension = FuncExtOpenCLCGenericAddressSpace in { + defm : VloadVstore<[GenericAS], 1>; } -let MinVersion = CL20 in { - foreach AS = [GenericAS] in { - def : Builtin<"vload_half", [Float, Size, PointerType<ConstType<Half>, AS>]>; - def : Builtin<"vloada_half", [Float, Size, PointerType<ConstType<Half>, AS>]>; +// 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] in { - def : Builtin<name, [VectorType<Float, VSize>, Size, PointerType<ConstType<Half>, AS>]>; + foreach name = ["vload_half" # VSize, "vloada_half" # VSize] in { + def : Builtin<name, [VectorType<Float, VSize>, Size, !cast<Type>("HalfPtrConst" # AS)], Attr.Pure>; } } - 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>]>; + if defStores then { + foreach rnd = ["", "_rte", "_rtz", "_rtp", "_rtn"] in { + foreach name = ["vstore_half" # rnd] in { + 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, !cast<Type>("HalfPtr" # AS)]>; + def : Builtin<name, [Void, VectorType<Double, VSize>, Size, !cast<Type>("HalfPtr" # 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>]>; - } - } +let Extension = FuncExtOpenCLCNamedAddressSpaceBuiltins in { + defm : VloadVstoreHalf<[GlobalAS, LocalAS, PrivateAS], 1>; +} +let Extension = FuncExtOpenCLCGenericAddressSpace in { + defm : VloadVstoreHalf<[GenericAS], 1>; } +// vload_half and vloada_half with constant address space are available regardless of version. +defm : VloadVstoreHalf<[ConstantAS], 0>; // OpenCL v3.0 s6.15.8 - Synchronization Functions. def : Builtin<"barrier", [Void, MemFenceFlags], Attr.Convergent>; @@ -958,7 +938,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. @@ -1099,42 +1079,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]>; + } + } } } @@ -1142,39 +1141,71 @@ 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>; + + 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 { + 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"); - 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_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>; + } + } } //-------------------------------------------------------------------- @@ -1241,30 +1272,35 @@ foreach coordTy = [Int, Float] in { } // --- 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 { - 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>; + defm : ImageReadSamplerless<"RO">; + let Extension = FuncExtOpenCLCReadWriteImages in { + defm : ImageReadSamplerless<"RW">; } } // --- 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>]>; @@ -1294,8 +1330,13 @@ foreach aQual = ["WO", "RW"] in { 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 { @@ -1319,6 +1360,12 @@ foreach aQual = ["RO", "WO", "RW"] in { } } +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 { @@ -1339,7 +1386,7 @@ 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>; @@ -1352,10 +1399,14 @@ let MinVersion = CL12 in { } } } + 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>]>; @@ -1366,11 +1417,17 @@ foreach aQual = ["WO", "RW"] in { } } +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>; } @@ -1395,7 +1452,9 @@ let MinVersion = CL20 in { // --- 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 @@ -1410,7 +1469,7 @@ def : Builtin<"is_valid_reserve_id", [Bool, ReserveId]>; // 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>]>; @@ -1537,14 +1596,21 @@ let Extension = FuncExtKhrMipmapImage in { } } } - // Added to section 6.13.14.5 - foreach aQual = ["RO", "WO", "RW"] in { - foreach imgTy = [Image1d, Image2d, Image3d, Image1dArray, Image2dArray, Image2dDepth, Image2dArrayDepth] in { - def : Builtin<"get_image_num_mip_levels", [Int, ImageType<imgTy, aQual>]>; - } +} + +// Added to section 6.13.14.5 +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. @@ -1571,53 +1637,57 @@ 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>]>; } } } //-------------------------------------------------------------------- // OpenCL Extension v2.0 s18.3 - Creating OpenCL Memory Objects from OpenGL MSAA Textures -let Extension = FuncExtKhrGlMsaaSharing in { - // --- Table 6.13.14.3 --- - foreach aQual = ["RO", "RW"] in { - 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.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>; } - - // --- 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>; - } + foreach imgTy = [Image2dArrayMsaa, Image2dArrayMsaaDepth] in { + def : Builtin<"get_image_array_size", [Size, ImageType<imgTy, aQual>], Attr.Const>; } } +let Extension = FuncExtKhrGlMsaaSharing in { + defm : ImageReadMsaa<"RO">; + defm : ImageQueryMsaa<"RO">; + defm : ImageQueryMsaa<"WO">; + defm : ImageReadMsaa<"RW">; + defm : ImageQueryMsaa<"RW">; +} + //-------------------------------------------------------------------- // OpenCL Extension v2.0 s28 - Subgroups // --- Table 28.2.1 --- @@ -1637,7 +1707,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 --- @@ -1774,6 +1846,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 { |