From 79a222fcf8c40a8a90e76118e2412e8697d02a89 Mon Sep 17 00:00:00 2001 From: Sven van Haastregt Date: Mon, 3 Jun 2019 09:39:11 +0000 Subject: [PATCH] [OpenCL] Declare builtin functions using TableGen This patch adds a `-fdeclare-opencl-builtins` command line option to the clang frontend. This enables clang to verify OpenCL C builtin function declarations using a fast StringMatcher lookup, instead of including the opencl-c.h file with the `-finclude-default-header` option. This avoids the large parse time penalty of the header file. This commit only adds the basic infrastructure and some of the OpenCL builtins. It does not cover all builtins defined by the various OpenCL specifications. As such, it is not a replacement for `-finclude-default-header` yet. RFC: http://lists.llvm.org/pipermail/cfe-dev/2018-November/060041.html Co-authored-by: Pierre Gondois Co-authored-by: Joey Gouly Co-authored-by: Sven van Haastregt Differential Revision: https://reviews.llvm.org/D60763 llvm-svn: 362371 --- clang/include/clang/Basic/CMakeLists.txt | 6 + clang/include/clang/Basic/LangOptions.def | 1 + clang/include/clang/Basic/OpenCLBuiltins.td | 296 +++++++++++++++++++ clang/include/clang/Driver/CC1Options.td | 4 +- clang/lib/Frontend/CompilerInvocation.cpp | 3 +- clang/lib/Sema/SemaLookup.cpp | 84 ++++++ clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl | 24 ++ clang/utils/TableGen/CMakeLists.txt | 1 + clang/utils/TableGen/ClangOpenCLBuiltinEmitter.cpp | 318 +++++++++++++++++++++ clang/utils/TableGen/TableGen.cpp | 6 + clang/utils/TableGen/TableGenBackends.h | 3 + 11 files changed, 744 insertions(+), 2 deletions(-) create mode 100644 clang/include/clang/Basic/OpenCLBuiltins.td create mode 100644 clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl create mode 100644 clang/utils/TableGen/ClangOpenCLBuiltinEmitter.cpp diff --git a/clang/include/clang/Basic/CMakeLists.txt b/clang/include/clang/Basic/CMakeLists.txt index 15bed5a..e26e683 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 3590757..ae3aeb4 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 0000000..7e37e55 --- /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 = _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]>; + } +} diff --git a/clang/include/clang/Driver/CC1Options.td b/clang/include/clang/Driver/CC1Options.td index d2d4710..76b36a1 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 9658080..717278c 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 f1d2a05..7643a06 100644 --- a/clang/lib/Sema/SemaLookup.cpp +++ b/clang/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 @@ 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 , 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 @@ 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 0000000..a19664f --- /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 dba0c94..3fc87d6 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 0000000..1e49503 --- /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 . +// 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 diff --git a/clang/utils/TableGen/TableGen.cpp b/clang/utils/TableGen/TableGen.cpp index 351768f..b9ec90f 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 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 08edb68..02af66c 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); -- 2.7.4