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