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