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