summaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
Diffstat
-rw-r--r--clang/include/clang/Basic/CMakeLists.txt6
-rw-r--r--clang/include/clang/Basic/LangOptions.def1
-rw-r--r--clang/include/clang/Basic/OpenCLBuiltins.td296
-rw-r--r--clang/include/clang/Driver/CC1Options.td4
-rw-r--r--clang/lib/Frontend/CompilerInvocation.cpp3
-rw-r--r--clang/lib/Sema/SemaLookup.cpp84
-rw-r--r--clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl24
-rw-r--r--clang/utils/TableGen/CMakeLists.txt1
-rw-r--r--clang/utils/TableGen/ClangOpenCLBuiltinEmitter.cpp318
-rw-r--r--clang/utils/TableGen/TableGen.cpp6
-rw-r--r--clang/utils/TableGen/TableGenBackends.h3
11 files changed, 744 insertions, 2 deletions
diff --git a/clang/include/clang/Basic/CMakeLists.txt b/clang/include/clang/Basic/CMakeLists.txt
index 15bed5adec9..e26e683b9ab 100644
--- a/clang/include/clang/Basic/CMakeLists.txt
+++ b/clang/include/clang/Basic/CMakeLists.txt
@@ -41,6 +41,12 @@ clang_tablegen(AttrHasAttributeImpl.inc -gen-clang-attr-has-attribute-impl
TARGET ClangAttrHasAttributeImpl
)
+clang_tablegen(OpenCLBuiltins.inc
+ -I ${CMAKE_CURRENT_SOURCE_DIR}/../../ -gen-clang-opencl-builtins
+ SOURCE OpenCLBuiltins.td
+ TARGET ClangOpenCLBuiltinsImpl
+ )
+
# ARM NEON
clang_tablegen(arm_neon.inc -gen-arm-neon-sema
SOURCE arm_neon.td
diff --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def
index 359075717f5..ae3aeb47019 100644
--- a/clang/include/clang/Basic/LangOptions.def
+++ b/clang/include/clang/Basic/LangOptions.def
@@ -256,6 +256,7 @@ LANGOPT(CFProtectionBranch , 1, 0, "Control-Flow Branch Protection enabled")
LANGOPT(FakeAddressSpaceMap , 1, 0, "OpenCL fake address space map")
ENUM_LANGOPT(AddressSpaceMapMangling , AddrSpaceMapMangling, 2, ASMM_Target, "OpenCL address space map mangling mode")
LANGOPT(IncludeDefaultHeader, 1, 0, "Include default header file for OpenCL")
+LANGOPT(DeclareOpenCLBuiltins, 1, 0, "Declare OpenCL builtin functions")
BENIGN_LANGOPT(DelayedTemplateParsing , 1, 0, "delayed template parsing")
LANGOPT(BlocksRuntimeOptional , 1, 0, "optional blocks runtime")
LANGOPT(
diff --git a/clang/include/clang/Basic/OpenCLBuiltins.td b/clang/include/clang/Basic/OpenCLBuiltins.td
new file mode 100644
index 00000000000..7e37e55dbaf
--- /dev/null
+++ b/clang/include/clang/Basic/OpenCLBuiltins.td
@@ -0,0 +1,296 @@
+//==--- OpenCLBuiltins.td - OpenCL builtin declarations -------------------===//
+//
+// The LLVM Compiler Infrastructure
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file contains TableGen definitions for OpenCL builtin function
+// declarations. In case of an unresolved function name in OpenCL, Clang will
+// check for a function described in this file when -fdeclare-opencl-builtins
+// is specified.
+//
+//===----------------------------------------------------------------------===//
+
+//===----------------------------------------------------------------------===//
+// Definitions of miscellaneous basic entities.
+//===----------------------------------------------------------------------===//
+// Versions of OpenCL
+class Version<int _Version> {
+ int Version = _Version;
+}
+def CL10: Version<100>;
+def CL11: Version<110>;
+def CL12: Version<120>;
+def CL20: Version<200>;
+
+// Address spaces
+// Pointer types need to be assigned an address space.
+class AddressSpace<string _AS> {
+ string AddrSpace = _AS;
+}
+def default_as : AddressSpace<"clang::LangAS::Default">;
+def private_as : AddressSpace<"clang::LangAS::opencl_private">;
+def global_as : AddressSpace<"clang::LangAS::opencl_global">;
+def constant_as : AddressSpace<"clang::LangAS::opencl_constant">;
+def local_as : AddressSpace<"clang::LangAS::opencl_local">;
+def generic_as : AddressSpace<"clang::LangAS::opencl_generic">;
+
+
+// Qualified Type. Allow to retrieve one ASTContext QualType.
+class QualType<string _Name> {
+ // Name of the field or function in a clang::ASTContext
+ // E.g. Name="IntTy" for the int type, and "getIntPtrType()" for an intptr_t
+ string Name = _Name;
+}
+
+// Helper class to store type access qualifiers (volatile, const, ...).
+class Qualifier<string _QualName> {
+ string QualName = _QualName;
+}
+
+//===----------------------------------------------------------------------===//
+// OpenCL C classes for types
+//===----------------------------------------------------------------------===//
+// OpenCL types (int, float, ...)
+class Type<string _Name, QualType _QTName> {
+ // Name of the Type
+ string Name = _Name;
+ // QualType associated with this type
+ QualType QTName = _QTName;
+ // Size of the vector (if applicable)
+ int VecWidth = 0;
+ // Is pointer
+ bit IsPointer = 0;
+ // List of qualifiers associated with the type (volatile, ...)
+ list<Qualifier> QualList = [];
+ // Address space
+ string AddrSpace = "clang::LangAS::Default";
+ // Access qualifier. Must be one of ("RO", "WO", "RW").
+ string AccessQualifier = "";
+}
+
+// OpenCL vector types (e.g. int2, int3, int16, float8, ...)
+class VectorType<Type _Ty, int _VecWidth> : Type<_Ty.Name, _Ty.QTName> {
+ int VecWidth = _VecWidth;
+}
+
+// OpenCL pointer types (e.g. int*, float*, ...)
+class PointerType<Type _Ty, AddressSpace _AS = global_as> :
+ Type<_Ty.Name, _Ty.QTName> {
+ bit IsPointer = 1;
+ string AddrSpace = _AS.AddrSpace;
+}
+
+// OpenCL image types (e.g. image2d_t, ...)
+class ImageType<Type _Ty, QualType _QTName, string _AccessQualifier> :
+ Type<_Ty.Name, _QTName> {
+ let AccessQualifier = _AccessQualifier;
+}
+
+//===----------------------------------------------------------------------===//
+// OpenCL C class for builtin functions
+//===----------------------------------------------------------------------===//
+class Builtin<string _Name, list<Type> _Signature> {
+ // Name of the builtin function
+ string Name = _Name;
+ // List of types used by the function. The first one is the return type and
+ // the following are the arguments. The list must have at least one element
+ // (the return type).
+ list<Type> Signature = _Signature;
+ // OpenCL Extension to which the function belongs (cl_khr_subgroups, ...)
+ string Extension = "";
+ // OpenCL Version to which the function belongs (CL10, ...)
+ Version Version = CL10;
+}
+
+//===----------------------------------------------------------------------===//
+// Multiclass definitions
+//===----------------------------------------------------------------------===//
+// multiclass BifN: Creates Builtin class instances for OpenCL builtin
+// functions with N arguments.
+// _Name : Name of the function
+// _Signature : Signature of the function (list of the Type used by the
+// function, the first one being the return type).
+// _IsVector : List of bit indicating if the type in the _Signature at the
+// same index is to be a vector in the multiple overloads. The
+// list must have at least one non-zero value.
+multiclass Bif0<string _Name, list<Type> _Signature, list<bit> _IsVector> {
+ def : Builtin<_Name, _Signature>;
+ foreach v = [2, 3, 4, 8, 16] in {
+ def : Builtin<_Name,
+ [!if(_IsVector[0], VectorType<_Signature[0], v>, _Signature[0])]>;
+ }
+}
+multiclass Bif1<string _Name, list<Type> _Signature, list<bit> _IsVector> {
+ def : Builtin<_Name, _Signature>;
+ foreach v = [2, 3, 4, 8, 16] in {
+ def : Builtin<_Name,
+ [!if(_IsVector[0], VectorType<_Signature[0], v>, _Signature[0]),
+ !if(_IsVector[1], VectorType<_Signature[1], v>, _Signature[1])]>;
+ }
+}
+multiclass Bif2<string _Name, list<Type> _Signature, list<bit> _IsVector> {
+ def : Builtin<_Name, _Signature>;
+ foreach v = [2, 3, 4, 8, 16] in {
+ def : Builtin<_Name,
+ [!if(_IsVector[0], VectorType<_Signature[0], v>, _Signature[0]),
+ !if(_IsVector[1], VectorType<_Signature[1], v>, _Signature[1]),
+ !if(_IsVector[2], VectorType<_Signature[2], v>, _Signature[2])]>;
+ }
+}
+multiclass Bif3<string _Name, list<Type> _Signature, list<bit> _IsVector> {
+ def : Builtin<_Name, _Signature>;
+ foreach v = [2, 3, 4, 8, 16] in {
+ def : Builtin<_Name,
+ [!if(_IsVector[0], VectorType<_Signature[0], v>, _Signature[0]),
+ !if(_IsVector[1], VectorType<_Signature[1], v>, _Signature[1]),
+ !if(_IsVector[2], VectorType<_Signature[2], v>, _Signature[2]),
+ !if(_IsVector[3], VectorType<_Signature[3], v>, _Signature[3])]>;
+ }
+}
+//===----------------------------------------------------------------------===//
+// Definitions of OpenCL C types
+//===----------------------------------------------------------------------===//
+// OpenCL v1.2 s6.1.1: Built-in Scalar Data Types
+def bool_t : Type<"bool", QualType<"BoolTy">>;
+def char_t : Type<"char", QualType<"CharTy">>;
+def uchar_t : Type<"uchar", QualType<"UnsignedCharTy">>;
+def short_t : Type<"short", QualType<"ShortTy">>;
+def ushort_t : Type<"ushort", QualType<"UnsignedShortTy">>;
+def int_t : Type<"int", QualType<"IntTy">>;
+def uint_t : Type<"uint", QualType<"UnsignedIntTy">>;
+def long_t : Type<"long", QualType<"LongTy">>;
+def ulong_t : Type<"ulong", QualType<"UnsignedLongTy">>;
+def float_t : Type<"float", QualType<"FloatTy">>;
+def double_t : Type<"double", QualType<"DoubleTy">>;
+def half_t : Type<"half", QualType<"HalfTy">>;
+def size_t : Type<"size_t", QualType<"getSizeType()">>;
+def ptrdiff_t : Type<"ptrdiff_t", QualType<"getPointerDiffType()">>;
+def intptr_t : Type<"intptr_t", QualType<"getIntPtrType()">>;
+def uintptr_t : Type<"uintptr_t", QualType<"getUIntPtrType()">>;
+def void_t : Type<"void", QualType<"VoidTy">>;
+
+// OpenCL v1.2 s6.1.2: Built-in Vector Data Types
+foreach v = [2, 3, 4, 8, 16] in {
+ def char#v#_t : VectorType<char_t, v>;
+ def uchar#v#_t : VectorType<uchar_t, v>;
+ def short#v#_t : VectorType<short_t, v>;
+ def ushort#v#_t : VectorType<ushort_t, v>;
+ def "int"#v#_t : VectorType<int_t, v>;
+ def uint#v#_t : VectorType<uint_t, v>;
+ def long#v#_t : VectorType<long_t, v>;
+ def ulong#v#_t : VectorType<ulong_t, v>;
+ def float#v#_t : VectorType<float_t, v>;
+ def double#v#_t : VectorType<double_t, v>;
+ def half#v#_t : VectorType<half_t, v>;
+}
+
+// OpenCL v1.2 s6.1.3: Other Built-in Data Types
+// These definitions with a "null" name are "abstract". They should not
+// be used in definitions of Builtin functions.
+def image2d_t : Type<"image2d_t", QualType<"null">>;
+def image3d_t : Type<"image3d_t", QualType<"null">>;
+def image2d_array_t : Type<"image2d_array_t", QualType<"null">>;
+def image1d_t : Type<"image1d_t", QualType<"null">>;
+def image1d_buffer_t : Type<"image1d_buffer_t", QualType<"null">>;
+def image1d_array_t : Type<"image1d_array_t", QualType<"null">>;
+// Unlike the few functions above, the following definitions can be used
+// in definitions of Builtin functions (they have a QualType with a name).
+foreach v = ["RO", "WO", "RW"] in {
+ def image2d_#v#_t : ImageType<image2d_t,
+ QualType<"OCLImage2d"#v#"Ty">,
+ v>;
+ def image3d_#v#_t : ImageType<image3d_t,
+ QualType<"OCLImage3d"#v#"Ty">,
+ v>;
+ def image2d_array#v#_t : ImageType<image2d_array_t,
+ QualType<"OCLImage2dArray"#v#"Ty">,
+ v>;
+ def image1d_#v#_t : ImageType<image1d_t,
+ QualType<"OCLImage1d"#v#"Ty">,
+ v>;
+ def image1d_buffer#v#_t : ImageType<image1d_buffer_t,
+ QualType<"OCLImage1dBuffer"#v#"Ty">,
+ v>;
+ def image1d_array#v#_t : ImageType<image1d_array_t,
+ QualType<"OCLImage1dArray"#v#"Ty">,
+ v>;
+}
+
+def sampler_t : Type<"sampler_t", QualType<"OCLSamplerTy">>;
+def event_t : Type<"event_t", QualType<"OCLEventTy">>;
+
+//===----------------------------------------------------------------------===//
+// Definitions of OpenCL builtin functions
+//===----------------------------------------------------------------------===//
+// OpenCL v1.2 s6.2.3: Explicit Conversions
+// Generate the convert_ builtins.
+foreach RType = [float_t, double_t, char_t, uchar_t, short_t, ushort_t,
+ int_t, uint_t, long_t, ulong_t] in {
+ foreach IType = [float_t, double_t, char_t, uchar_t, short_t, ushort_t,
+ int_t, uint_t, long_t, ulong_t] in {
+ foreach sat = ["", "_sat"] in {
+ foreach rte = ["", "_rte", "_rtz", "_rtp", "_rtn"] in {
+ def : Builtin<"convert_" # RType.Name # sat # rte, [RType, IType]>;
+ foreach v = [2, 3, 4, 8, 16] in {
+ def : Builtin<"convert_" # RType.Name # v # sat # rte,
+ [VectorType<RType, v>,
+ VectorType<IType, v>]>;
+ }
+ }
+ }
+ }
+}
+
+// OpenCL v1.2 s6.12.1: Work-Item Functions
+def get_work_dim : Builtin<"get_work_dim", [uint_t]>;
+foreach name = ["get_global_size", "get_global_id", "get_local_size",
+ "get_local_id", "get_num_groups", "get_group_id",
+ "get_global_offset"] in {
+ def : Builtin<name, [size_t, uint_t]>;
+}
+
+// OpenCL v1.2 s6.12.2: Math Functions
+foreach name = ["acos", "acosh", "acospi",
+ "asin", "asinh", "asinpi",
+ "atan", "atanh", "atanpi"] in {
+ foreach type = [float_t, double_t, half_t] in {
+ defm : Bif1<name, [type, type], [1, 1]>;
+ }
+}
+
+foreach name = ["atan2", "atan2pi"] in {
+ foreach type = [float_t, double_t, half_t] in {
+ defm : Bif2<name, [type, type, type], [1, 1, 1]>;
+ }
+}
+
+foreach name = ["fmax", "fmin"] in {
+ foreach type = [float_t, double_t, half_t] in {
+ defm : Bif2<name, [type, type, type], [1, 1, 1]>;
+ defm : Bif2<name, [type, type, type], [1, 1, 0]>;
+ }
+}
+
+// OpenCL v1.2 s6.12.14: Built-in Image Read Functions
+def read_imagef : Builtin<"read_imagef",
+ [float4_t, image2d_RO_t, VectorType<int_t, 2>]>;
+def write_imagef : Builtin<"write_imagef",
+ [void_t,
+ image2d_WO_t,
+ VectorType<int_t, 2>,
+ VectorType<float_t, 4>]>;
+
+
+// OpenCL v2.0 s9.17.3: Additions to section 6.13.1: Work-Item Functions
+let Version = CL20 in {
+ let Extension = "cl_khr_subgroups" in {
+ def get_sub_group_size : Builtin<"get_sub_group_size", [uint_t]>;
+ def get_max_sub_group_size : Builtin<"get_max_sub_group_size", [uint_t]>;
+ def get_num_sub_groups : Builtin<"get_num_sub_groups", [uint_t]>;
+ }
+}
diff --git a/clang/include/clang/Driver/CC1Options.td b/clang/include/clang/Driver/CC1Options.td
index d2d471089e5..76b36a18269 100644
--- a/clang/include/clang/Driver/CC1Options.td
+++ b/clang/include/clang/Driver/CC1Options.td
@@ -778,7 +778,9 @@ def fallow_half_arguments_and_returns : Flag<["-"], "fallow-half-arguments-and-r
def fdefault_calling_conv_EQ : Joined<["-"], "fdefault-calling-conv=">,
HelpText<"Set default calling convention">, Values<"cdecl,fastcall,stdcall,vectorcall,regcall">;
def finclude_default_header : Flag<["-"], "finclude-default-header">,
- HelpText<"Include the default header file for OpenCL">;
+ HelpText<"Include default header file for OpenCL">;
+def fdeclare_opencl_builtins : Flag<["-"], "fdeclare-opencl-builtins">,
+ HelpText<"Add OpenCL builtin function declarations (experimental)">;
def fpreserve_vec3_type : Flag<["-"], "fpreserve-vec3-type">,
HelpText<"Preserve 3-component vector type">;
def fwchar_type_EQ : Joined<["-"], "fwchar-type=">,
diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp
index 96580804576..717278c0861 100644
--- a/clang/lib/Frontend/CompilerInvocation.cpp
+++ b/clang/lib/Frontend/CompilerInvocation.cpp
@@ -2179,7 +2179,7 @@ void CompilerInvocation::setLangDefaults(LangOptions &Opts, InputKind IK,
Opts.NativeHalfArgsAndReturns = 1;
Opts.OpenCLCPlusPlus = Opts.CPlusPlus;
// Include default header file for OpenCL.
- if (Opts.IncludeDefaultHeader) {
+ if (Opts.IncludeDefaultHeader && !Opts.DeclareOpenCLBuiltins) {
PPOpts.Includes.push_back("opencl-c.h");
}
}
@@ -2385,6 +2385,7 @@ static void ParseLangArgs(LangOptions &Opts, ArgList &Args, InputKind IK,
}
Opts.IncludeDefaultHeader = Args.hasArg(OPT_finclude_default_header);
+ Opts.DeclareOpenCLBuiltins = Args.hasArg(OPT_fdeclare_opencl_builtins);
llvm::Triple T(TargetOpts.Triple);
CompilerInvocation::setLangDefaults(Opts, IK, T, PPOpts, LangStd);
diff --git a/clang/lib/Sema/SemaLookup.cpp b/clang/lib/Sema/SemaLookup.cpp
index f1d2a052404..7643a06a82f 100644
--- a/clang/lib/Sema/SemaLookup.cpp
+++ b/clang/lib/Sema/SemaLookup.cpp
@@ -46,6 +46,8 @@
#include <utility>
#include <vector>
+#include "clang/Basic/OpenCLBuiltins.inc"
+
using namespace clang;
using namespace sema;
@@ -670,6 +672,79 @@ LLVM_DUMP_METHOD void LookupResult::dump() {
D->dump();
}
+/// When trying to resolve a function name, if the isOpenCLBuiltin function
+/// defined in "OpenCLBuiltins.inc" returns a non-null <Index, Len>, then the
+/// identifier is referencing an OpenCL builtin function. Thus, all its
+/// prototypes are added to the LookUpResult.
+///
+/// \param S The Sema instance
+/// \param LR The LookupResult instance
+/// \param II The identifier being resolved
+/// \param Index The list of prototypes starts at Index in OpenCLBuiltins[]
+/// \param Len The list of prototypes has Len elements
+static void InsertOCLBuiltinDeclarations(Sema &S, LookupResult &LR,
+ IdentifierInfo *II, unsigned Index,
+ unsigned Len) {
+
+ for (unsigned i = 0; i < Len; ++i) {
+ OpenCLBuiltinDecl &Decl = OpenCLBuiltins[Index - 1 + i];
+ ASTContext &Context = S.Context;
+
+ // Ignore this BIF if the version is incorrect.
+ if (Context.getLangOpts().OpenCLVersion < Decl.Version)
+ continue;
+
+ FunctionProtoType::ExtProtoInfo PI;
+ PI.Variadic = false;
+
+ // Defined in "OpenCLBuiltins.inc"
+ QualType RT = OCL2Qual(Context, OpenCLSignature[Decl.ArgTableIndex]);
+
+ SmallVector<QualType, 5> ArgTypes;
+ for (unsigned I = 1; I < Decl.NumArgs; I++) {
+ QualType Ty = OCL2Qual(Context, OpenCLSignature[Decl.ArgTableIndex + I]);
+ ArgTypes.push_back(Ty);
+ }
+
+ QualType R = Context.getFunctionType(RT, ArgTypes, PI);
+ SourceLocation Loc = LR.getNameLoc();
+
+ // TODO: This part is taken from Sema::LazilyCreateBuiltin,
+ // maybe refactor it.
+ DeclContext *Parent = Context.getTranslationUnitDecl();
+ FunctionDecl *New = FunctionDecl::Create(Context, Parent, Loc, Loc, II, R,
+ /*TInfo=*/nullptr, SC_Extern,
+ false, R->isFunctionProtoType());
+ New->setImplicit();
+
+ // Create Decl objects for each parameter, adding them to the
+ // FunctionDecl.
+ if (const FunctionProtoType *FT = dyn_cast<FunctionProtoType>(R)) {
+ SmallVector<ParmVarDecl *, 16> Params;
+ for (unsigned i = 0, e = FT->getNumParams(); i != e; ++i) {
+ ParmVarDecl *Parm =
+ ParmVarDecl::Create(Context, New, SourceLocation(),
+ SourceLocation(), nullptr, FT->getParamType(i),
+ /*TInfo=*/nullptr, SC_None, nullptr);
+ Parm->setScopeInfo(0, i);
+ Params.push_back(Parm);
+ }
+ New->setParams(Params);
+ }
+
+ New->addAttr(OverloadableAttr::CreateImplicit(Context));
+
+ if (strlen(Decl.Extension))
+ S.setOpenCLExtensionForDecl(New, Decl.Extension);
+
+ LR.addDecl(New);
+ }
+
+ // If we added overloads, need to resolve the lookup result.
+ if (Len > 1)
+ LR.resolveKind();
+}
+
/// Lookup a builtin function, when name lookup would otherwise
/// fail.
static bool LookupBuiltin(Sema &S, LookupResult &R) {
@@ -692,6 +767,15 @@ static bool LookupBuiltin(Sema &S, LookupResult &R) {
}
}
+ // Check if this is an OpenCL Builtin, and if so, insert its overloads.
+ if (S.getLangOpts().OpenCL && S.getLangOpts().DeclareOpenCLBuiltins) {
+ auto Index = isOpenCLBuiltin(II->getName());
+ if (Index.first) {
+ InsertOCLBuiltinDeclarations(S, R, II, Index.first, Index.second);
+ return true;
+ }
+ }
+
// If this is a builtin on this (or all) targets, create the decl.
if (unsigned BuiltinID = II->getBuiltinID()) {
// In C++ and OpenCL (spec v1.2 s6.9.f), we don't have any predefined
diff --git a/clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl b/clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl
new file mode 100644
index 00000000000..a19664f3104
--- /dev/null
+++ b/clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl
@@ -0,0 +1,24 @@
+// RUN: %clang_cc1 %s -triple spir -verify -pedantic -fsyntax-only -cl-std=CL2.0 -fdeclare-opencl-builtins
+
+// Test the -fdeclare-opencl-builtins option.
+
+typedef float float4 __attribute__((ext_vector_type(4)));
+typedef int int4 __attribute__((ext_vector_type(4)));
+typedef int int2 __attribute__((ext_vector_type(2)));
+typedef unsigned int uint;
+typedef __SIZE_TYPE__ size_t;
+
+kernel void basic_conversion(global float4 *buf, global int4 *res) {
+ res[0] = convert_int4(buf[0]);
+}
+
+kernel void basic_readonly_image_type(__read_only image2d_t img, int2 coord, global float4 *out) {
+ out[0] = read_imagef(img, coord);
+}
+
+kernel void basic_subgroup(global uint *out) {
+ out[0] = get_sub_group_size();
+// expected-error@-1{{use of declaration 'get_sub_group_size' requires cl_khr_subgroups extension to be enabled}}
+#pragma OPENCL EXTENSION cl_khr_subgroups : enable
+ out[1] = get_sub_group_size();
+}
diff --git a/clang/utils/TableGen/CMakeLists.txt b/clang/utils/TableGen/CMakeLists.txt
index dba0c94ac0e..3fc87d65525 100644
--- a/clang/utils/TableGen/CMakeLists.txt
+++ b/clang/utils/TableGen/CMakeLists.txt
@@ -8,6 +8,7 @@ add_tablegen(clang-tblgen CLANG
ClangCommentHTMLTagsEmitter.cpp
ClangDataCollectorsEmitter.cpp
ClangDiagnosticsEmitter.cpp
+ ClangOpenCLBuiltinEmitter.cpp
ClangOptionDocEmitter.cpp
ClangSACheckersEmitter.cpp
NeonEmitter.cpp
diff --git a/clang/utils/TableGen/ClangOpenCLBuiltinEmitter.cpp b/clang/utils/TableGen/ClangOpenCLBuiltinEmitter.cpp
new file mode 100644
index 00000000000..1e495039c49
--- /dev/null
+++ b/clang/utils/TableGen/ClangOpenCLBuiltinEmitter.cpp
@@ -0,0 +1,318 @@
+//===- ClangOpenCLBuiltinEmitter.cpp - Generate Clang OpenCL Builtin handling
+//
+// The LLVM Compiler Infrastructure
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This tablegen backend emits code for checking whether a function is an
+// OpenCL builtin function. If so, all overloads of this function are
+// added to the LookupResult. The generated include file is used by
+// SemaLookup.cpp
+//
+// For a successful lookup of e.g. the "cos" builtin, isOpenCLBuiltin("cos")
+// returns a pair <Index, Len>.
+// OpenCLBuiltins[Index] to OpenCLBuiltins[Index + Len] contains the pairs
+// <SigIndex, SigLen> of the overloads of "cos".
+// OpenCLSignature[SigIndex] to OpenCLSignature[SigIndex + SigLen] contains
+// one of the signatures of "cos". The OpenCLSignature entry can be
+// referenced by other functions, i.e. "sin", since multiple OpenCL builtins
+// share the same signature.
+//===----------------------------------------------------------------------===//
+
+#include "llvm/ADT/MapVector.h"
+#include "llvm/ADT/STLExtras.h"
+#include "llvm/ADT/SmallString.h"
+#include "llvm/ADT/StringExtras.h"
+#include "llvm/ADT/StringRef.h"
+#include "llvm/ADT/StringSet.h"
+#include "llvm/Support/ErrorHandling.h"
+#include "llvm/Support/raw_ostream.h"
+#include "llvm/TableGen/Error.h"
+#include "llvm/TableGen/Record.h"
+#include "llvm/TableGen/StringMatcher.h"
+#include "llvm/TableGen/TableGenBackend.h"
+#include <set>
+
+using namespace llvm;
+
+namespace {
+class BuiltinNameEmitter {
+public:
+ BuiltinNameEmitter(RecordKeeper &Records, raw_ostream &OS)
+ : Records(Records), OS(OS) {}
+
+ // Entrypoint to generate the functions and structures for checking
+ // whether a function is an OpenCL builtin function.
+ void Emit();
+
+private:
+ // Contains OpenCL builtin functions and related information, stored as
+ // Record instances. They are coming from the associated TableGen file.
+ RecordKeeper &Records;
+
+ // The output file.
+ raw_ostream &OS;
+
+ // Emit the enums and structs.
+ void EmitDeclarations();
+
+ // Parse the Records generated by TableGen and populate OverloadInfo and
+ // SignatureSet.
+ void GetOverloads();
+
+ // Emit the OpenCLSignature table. This table contains all possible
+ // signatures, and is a struct OpenCLType. A signature is composed of a
+ // return type (mandatory), followed by zero or more argument types.
+ // E.g.:
+ // // 12
+ // { OCLT_uchar, 4, clang::LangAS::Default, false },
+ // { OCLT_float, 4, clang::LangAS::Default, false },
+ // This means that index 12 represents a signature
+ // - returning a uchar vector of 4 elements, and
+ // - taking as first argument a float vector of 4 elements.
+ void EmitSignatureTable();
+
+ // Emit the OpenCLBuiltins table. This table contains all overloads of
+ // each function, and is a struct OpenCLBuiltinDecl.
+ // E.g.:
+ // // acos
+ // { 2, 0, "", 100 },
+ // This means that the signature of this acos overload is defined in OpenCL
+ // version 1.0 (100) and does not belong to any extension (""). It has a
+ // 1 argument (+1 for the return type), stored at index 0 in the
+ // OpenCLSignature table.
+ void EmitBuiltinTable();
+
+ // Emit a StringMatcher function to check whether a function name is an
+ // OpenCL builtin function name.
+ void EmitStringMatcher();
+
+ // Emit a function returning the clang QualType instance associated with
+ // the TableGen Record Type.
+ void EmitQualTypeFinder();
+
+ // Contains a list of the available signatures, without the name of the
+ // function. Each pair consists of a signature and a cumulative index.
+ // E.g.: <<float, float>, 0>,
+ // <<float, int, int, 2>>,
+ // <<float>, 5>,
+ // ...
+ // <<double, double>, 35>.
+ std::vector<std::pair<std::vector<Record *>, unsigned>> SignatureSet;
+
+ // Map the name of a builtin function to its prototypes (instances of the
+ // TableGen "Builtin" class).
+ // Each prototype is registered as a pair of:
+ // <pointer to the "Builtin" instance,
+ // cumulative index of the associated signature in the SignatureSet>
+ // E.g.: The function cos: (float cos(float), double cos(double), ...)
+ // <"cos", <<ptrToPrototype0, 5>,
+ // <ptrToPrototype1, 35>>
+ // <ptrToPrototype2, 79>>
+ // ptrToPrototype1 has the following signature: <double, double>
+ MapVector<StringRef, std::vector<std::pair<const Record *, unsigned>>>
+ OverloadInfo;
+};
+} // namespace
+
+void BuiltinNameEmitter::Emit() {
+ emitSourceFileHeader("OpenCL Builtin handling", OS);
+
+ OS << "#include \"llvm/ADT/StringRef.h\"\n";
+ OS << "using namespace clang;\n\n";
+
+ EmitDeclarations();
+
+ GetOverloads();
+
+ EmitSignatureTable();
+
+ EmitBuiltinTable();
+
+ EmitStringMatcher();
+
+ EmitQualTypeFinder();
+}
+
+void BuiltinNameEmitter::EmitDeclarations() {
+ OS << "enum OpenCLTypeID {\n";
+ std::vector<Record *> Types = Records.getAllDerivedDefinitions("Type");
+ StringMap<bool> TypesSeen;
+ for (const auto *T : Types) {
+ if (TypesSeen.find(T->getValueAsString("Name")) == TypesSeen.end())
+ OS << " OCLT_" + T->getValueAsString("Name") << ",\n";
+ TypesSeen.insert(std::make_pair(T->getValueAsString("Name"), true));
+ }
+ OS << "};\n";
+
+ OS << R"(
+
+// Type used in a prototype of an OpenCL builtin function.
+struct OpenCLType {
+ // A type (e.g.: float, int, ...)
+ OpenCLTypeID ID;
+ // Size of vector (if applicable)
+ unsigned VectorWidth;
+ // Address space of the pointer (if applicable)
+ LangAS AS;
+ // Whether the type is a pointer
+ bool isPointer;
+};
+
+// One overload of an OpenCL builtin function.
+struct OpenCLBuiltinDecl {
+ // Number of arguments for the signature
+ unsigned NumArgs;
+ // Index in the OpenCLSignature table to get the required types
+ unsigned ArgTableIndex;
+ // Extension to which it belongs (e.g. cl_khr_subgroups)
+ const char *Extension;
+ // Version in which it was introduced (e.g. CL20)
+ unsigned Version;
+};
+
+)";
+}
+
+void BuiltinNameEmitter::GetOverloads() {
+ unsigned CumulativeSignIndex = 0;
+ std::vector<Record *> Builtins = Records.getAllDerivedDefinitions("Builtin");
+ for (const auto *B : Builtins) {
+ StringRef BName = B->getValueAsString("Name");
+ if (OverloadInfo.find(BName) == OverloadInfo.end()) {
+ OverloadInfo.insert(std::make_pair(
+ BName, std::vector<std::pair<const Record *, unsigned>>{}));
+ }
+
+ auto Signature = B->getValueAsListOfDefs("Signature");
+ auto it =
+ std::find_if(SignatureSet.begin(), SignatureSet.end(),
+ [&](const std::pair<std::vector<Record *>, unsigned> &a) {
+ return a.first == Signature;
+ });
+ unsigned SignIndex;
+ if (it == SignatureSet.end()) {
+ SignatureSet.push_back(std::make_pair(Signature, CumulativeSignIndex));
+ SignIndex = CumulativeSignIndex;
+ CumulativeSignIndex += Signature.size();
+ } else {
+ SignIndex = it->second;
+ }
+ OverloadInfo[BName].push_back(std::make_pair(B, SignIndex));
+ }
+}
+
+void BuiltinNameEmitter::EmitSignatureTable() {
+ OS << "OpenCLType OpenCLSignature[] = {\n";
+ for (auto &P : SignatureSet) {
+ OS << "// " << P.second << "\n";
+ for (Record *R : P.first) {
+ OS << "{ OCLT_" << R->getValueAsString("Name") << ", "
+ << R->getValueAsInt("VecWidth") << ", "
+ << R->getValueAsString("AddrSpace") << ", "
+ << R->getValueAsBit("IsPointer") << "},";
+ OS << "\n";
+ }
+ }
+ OS << "};\n\n";
+}
+
+void BuiltinNameEmitter::EmitBuiltinTable() {
+ OS << "OpenCLBuiltinDecl OpenCLBuiltins[] = {\n";
+ for (auto &i : OverloadInfo) {
+ StringRef Name = i.first;
+ OS << "// " << Name << "\n";
+ for (auto &Overload : i.second) {
+ OS << " { " << Overload.first->getValueAsListOfDefs("Signature").size()
+ << ", " << Overload.second << ", " << '"'
+ << Overload.first->getValueAsString("Extension") << "\", "
+ << Overload.first->getValueAsDef("Version")->getValueAsInt("Version")
+ << " },\n";
+ }
+ }
+ OS << "};\n\n";
+}
+
+void BuiltinNameEmitter::EmitStringMatcher() {
+ std::vector<StringMatcher::StringPair> ValidBuiltins;
+ unsigned CumulativeIndex = 1;
+ for (auto &i : OverloadInfo) {
+ auto &Ov = i.second;
+ std::string RetStmt;
+ raw_string_ostream SS(RetStmt);
+ SS << "return std::make_pair(" << CumulativeIndex << ", " << Ov.size()
+ << ");";
+ SS.flush();
+ CumulativeIndex += Ov.size();
+
+ ValidBuiltins.push_back(StringMatcher::StringPair(i.first, RetStmt));
+ }
+
+ OS << R"(
+// Return 0 if name is not a recognized OpenCL builtin, or an index
+// into a table of declarations if it is an OpenCL builtin.
+std::pair<unsigned, unsigned> isOpenCLBuiltin(llvm::StringRef name) {
+
+)";
+
+ StringMatcher("name", ValidBuiltins, OS).Emit(0, true);
+
+ OS << " return std::make_pair(0, 0);\n";
+ OS << "}\n";
+}
+
+void BuiltinNameEmitter::EmitQualTypeFinder() {
+ OS << R"(
+
+static QualType OCL2Qual(ASTContext &Context, OpenCLType Ty) {
+ QualType RT = Context.VoidTy;
+ switch (Ty.ID) {
+)";
+
+ std::vector<Record *> Types = Records.getAllDerivedDefinitions("Type");
+ StringMap<bool> TypesSeen;
+
+ for (const auto *T : Types) {
+ // Check we have not seen this Type
+ if (TypesSeen.find(T->getValueAsString("Name")) != TypesSeen.end())
+ continue;
+ TypesSeen.insert(std::make_pair(T->getValueAsString("Name"), true));
+
+ // Check the Type does not have an "abstract" QualType
+ auto QT = T->getValueAsDef("QTName");
+ if (QT->getValueAsString("Name") == "null")
+ continue;
+
+ OS << " case OCLT_" << T->getValueAsString("Name") << ":\n";
+ OS << " RT = Context." << QT->getValueAsString("Name") << ";\n";
+ OS << " break;\n";
+ }
+ OS << " }\n";
+
+ // Special cases
+ OS << R"(
+ if (Ty.VectorWidth > 0)
+ RT = Context.getExtVectorType(RT, Ty.VectorWidth);
+
+ if (Ty.isPointer) {
+ RT = Context.getAddrSpaceQualType(RT, Ty.AS);
+ RT = Context.getPointerType(RT);
+ }
+
+ return RT;
+}
+)";
+}
+
+namespace clang {
+
+void EmitClangOpenCLBuiltins(RecordKeeper &Records, raw_ostream &OS) {
+ BuiltinNameEmitter NameChecker(Records, OS);
+ NameChecker.Emit();
+}
+
+} // end namespace clang
diff --git a/clang/utils/TableGen/TableGen.cpp b/clang/utils/TableGen/TableGen.cpp
index 351768fe965..b9ec90fd5bc 100644
--- a/clang/utils/TableGen/TableGen.cpp
+++ b/clang/utils/TableGen/TableGen.cpp
@@ -53,6 +53,7 @@ enum ActionType {
GenClangCommentHTMLNamedCharacterReferences,
GenClangCommentCommandInfo,
GenClangCommentCommandList,
+ GenClangOpenCLBuiltins,
GenArmNeon,
GenArmFP16,
GenArmNeonSema,
@@ -147,6 +148,8 @@ cl::opt<ActionType> Action(
clEnumValN(GenClangCommentCommandList, "gen-clang-comment-command-list",
"Generate list of commands that are used in "
"documentation comments"),
+ clEnumValN(GenClangOpenCLBuiltins, "gen-clang-opencl-builtins",
+ "Generate OpenCL builtin declaration handlers"),
clEnumValN(GenArmNeon, "gen-arm-neon", "Generate arm_neon.h for clang"),
clEnumValN(GenArmFP16, "gen-arm-fp16", "Generate arm_fp16.h for clang"),
clEnumValN(GenArmNeonSema, "gen-arm-neon-sema",
@@ -266,6 +269,9 @@ bool ClangTableGenMain(raw_ostream &OS, RecordKeeper &Records) {
case GenClangCommentCommandList:
EmitClangCommentCommandList(Records, OS);
break;
+ case GenClangOpenCLBuiltins:
+ EmitClangOpenCLBuiltins(Records, OS);
+ break;
case GenArmNeon:
EmitNeon(Records, OS);
break;
diff --git a/clang/utils/TableGen/TableGenBackends.h b/clang/utils/TableGen/TableGenBackends.h
index 08edb685742..02af66c5bf8 100644
--- a/clang/utils/TableGen/TableGenBackends.h
+++ b/clang/utils/TableGen/TableGenBackends.h
@@ -90,6 +90,9 @@ void EmitClangAttrDocs(llvm::RecordKeeper &Records, llvm::raw_ostream &OS);
void EmitClangDiagDocs(llvm::RecordKeeper &Records, llvm::raw_ostream &OS);
void EmitClangOptDocs(llvm::RecordKeeper &Records, llvm::raw_ostream &OS);
+void EmitClangOpenCLBuiltins(llvm::RecordKeeper &Records,
+ llvm::raw_ostream &OS);
+
void EmitClangDataCollectors(llvm::RecordKeeper &Records,
llvm::raw_ostream &OS);
OpenPOWER on IntegriCloud