Index: cfe/trunk/include/clang/Basic/CMakeLists.txt =================================================================== --- cfe/trunk/include/clang/Basic/CMakeLists.txt +++ cfe/trunk/include/clang/Basic/CMakeLists.txt @@ -41,6 +41,12 @@ 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 Index: cfe/trunk/include/clang/Basic/LangOptions.def =================================================================== --- cfe/trunk/include/clang/Basic/LangOptions.def +++ cfe/trunk/include/clang/Basic/LangOptions.def @@ -256,6 +256,7 @@ 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( Index: cfe/trunk/include/clang/Basic/OpenCLBuiltins.td =================================================================== --- cfe/trunk/include/clang/Basic/OpenCLBuiltins.td +++ cfe/trunk/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 = _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 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 { + // 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 = _QualName; +} + +//===----------------------------------------------------------------------===// +// OpenCL C classes for types +//===----------------------------------------------------------------------===// +// OpenCL types (int, float, ...) +class Type { + // 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 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.Name, _Ty.QTName> { + int VecWidth = _VecWidth; +} + +// OpenCL pointer types (e.g. int*, float*, ...) +class PointerType : + Type<_Ty.Name, _Ty.QTName> { + bit IsPointer = 1; + string AddrSpace = _AS.AddrSpace; +} + +// OpenCL image types (e.g. image2d_t, ...) +class ImageType : + Type<_Ty.Name, _QTName> { + let AccessQualifier = _AccessQualifier; +} + +//===----------------------------------------------------------------------===// +// OpenCL C class for builtin functions +//===----------------------------------------------------------------------===// +class Builtin _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 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 _Signature, list _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 _Signature, list _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 _Signature, list _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 _Signature, list _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; + def uchar#v#_t : VectorType; + def short#v#_t : VectorType; + def ushort#v#_t : VectorType; + def "int"#v#_t : VectorType; + def uint#v#_t : VectorType; + def long#v#_t : VectorType; + def ulong#v#_t : VectorType; + def float#v#_t : VectorType; + def double#v#_t : VectorType; + def half#v#_t : VectorType; +} + +// 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, + v>; + def image3d_#v#_t : ImageType, + v>; + def image2d_array#v#_t : ImageType, + v>; + def image1d_#v#_t : ImageType, + v>; + def image1d_buffer#v#_t : ImageType, + v>; + def image1d_array#v#_t : ImageType, + 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, + VectorType]>; + } + } + } + } +} + +// 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; +} + +// 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; + } +} + +foreach name = ["atan2", "atan2pi"] in { + foreach type = [float_t, double_t, half_t] in { + defm : Bif2; + } +} + +foreach name = ["fmax", "fmin"] in { + foreach type = [float_t, double_t, half_t] in { + defm : Bif2; + defm : Bif2; + } +} + +// OpenCL v1.2 s6.12.14: Built-in Image Read Functions +def read_imagef : Builtin<"read_imagef", + [float4_t, image2d_RO_t, VectorType]>; +def write_imagef : Builtin<"write_imagef", + [void_t, + image2d_WO_t, + VectorType, + VectorType]>; + + +// 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]>; + } +} Index: cfe/trunk/include/clang/Driver/CC1Options.td =================================================================== --- cfe/trunk/include/clang/Driver/CC1Options.td +++ cfe/trunk/include/clang/Driver/CC1Options.td @@ -778,7 +778,9 @@ 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=">, Index: cfe/trunk/lib/Frontend/CompilerInvocation.cpp =================================================================== --- cfe/trunk/lib/Frontend/CompilerInvocation.cpp +++ cfe/trunk/lib/Frontend/CompilerInvocation.cpp @@ -2179,7 +2179,7 @@ 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 @@ } 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); Index: cfe/trunk/lib/Sema/SemaLookup.cpp =================================================================== --- cfe/trunk/lib/Sema/SemaLookup.cpp +++ cfe/trunk/lib/Sema/SemaLookup.cpp @@ -46,6 +46,8 @@ #include #include +#include "clang/Basic/OpenCLBuiltins.inc" + using namespace clang; using namespace sema; @@ -670,6 +672,79 @@ D->dump(); } +/// When trying to resolve a function name, if the isOpenCLBuiltin function +/// defined in "OpenCLBuiltins.inc" returns a non-null , 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 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(R)) { + SmallVector 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 @@ } } + // 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 Index: cfe/trunk/test/SemaOpenCL/fdeclare-opencl-builtins.cl =================================================================== --- cfe/trunk/test/SemaOpenCL/fdeclare-opencl-builtins.cl +++ cfe/trunk/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(); +} Index: cfe/trunk/utils/TableGen/CMakeLists.txt =================================================================== --- cfe/trunk/utils/TableGen/CMakeLists.txt +++ cfe/trunk/utils/TableGen/CMakeLists.txt @@ -8,6 +8,7 @@ ClangCommentHTMLTagsEmitter.cpp ClangDataCollectorsEmitter.cpp ClangDiagnosticsEmitter.cpp + ClangOpenCLBuiltinEmitter.cpp ClangOptionDocEmitter.cpp ClangSACheckersEmitter.cpp NeonEmitter.cpp Index: cfe/trunk/utils/TableGen/ClangOpenCLBuiltinEmitter.cpp =================================================================== --- cfe/trunk/utils/TableGen/ClangOpenCLBuiltinEmitter.cpp +++ cfe/trunk/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 . +// OpenCLBuiltins[Index] to OpenCLBuiltins[Index + Len] contains the pairs +// 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 + +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.: <, 0>, + // <>, + // <, 5>, + // ... + // <, 35>. + std::vector, 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: + // + // E.g.: The function cos: (float cos(float), double cos(double), ...) + // <"cos", <, + // > + // > + // ptrToPrototype1 has the following signature: + MapVector>> + 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 Types = Records.getAllDerivedDefinitions("Type"); + StringMap 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 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>{})); + } + + auto Signature = B->getValueAsListOfDefs("Signature"); + auto it = + std::find_if(SignatureSet.begin(), SignatureSet.end(), + [&](const std::pair, 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 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 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 Types = Records.getAllDerivedDefinitions("Type"); + StringMap 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 Index: cfe/trunk/utils/TableGen/TableGen.cpp =================================================================== --- cfe/trunk/utils/TableGen/TableGen.cpp +++ cfe/trunk/utils/TableGen/TableGen.cpp @@ -53,6 +53,7 @@ GenClangCommentHTMLNamedCharacterReferences, GenClangCommentCommandInfo, GenClangCommentCommandList, + GenClangOpenCLBuiltins, GenArmNeon, GenArmFP16, GenArmNeonSema, @@ -147,6 +148,8 @@ 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 @@ case GenClangCommentCommandList: EmitClangCommentCommandList(Records, OS); break; + case GenClangOpenCLBuiltins: + EmitClangOpenCLBuiltins(Records, OS); + break; case GenArmNeon: EmitNeon(Records, OS); break; Index: cfe/trunk/utils/TableGen/TableGenBackends.h =================================================================== --- cfe/trunk/utils/TableGen/TableGenBackends.h +++ cfe/trunk/utils/TableGen/TableGenBackends.h @@ -90,6 +90,9 @@ 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);