mirror of
https://github.com/capstone-engine/llvm-capstone.git
synced 2024-12-21 15:02:51 +00:00
[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
This commit is contained in:
Sven van Haastregt
parent
301f8fd632
commit
79a222fcf8
@ -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
|
||||
|
||||
@ -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(
|
||||
|
||||
296
clang/include/clang/Basic/OpenCLBuiltins.td
Normal file
296
clang/include/clang/Basic/OpenCLBuiltins.td
Normal file
@ -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]>;
|
||||
}
|
||||
}
|
||||
@ -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=">,
|
||||
|
||||
@ -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);
|
||||
|
||||
@ -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
|
||||
|
||||
24
clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl
Normal file
24
clang/test/SemaOpenCL/fdeclare-opencl-builtins.cl
Normal file
@ -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();
|
||||
}
|
||||
@ -8,6 +8,7 @@ add_tablegen(clang-tblgen CLANG
|
||||
ClangCommentHTMLTagsEmitter.cpp
|
||||
ClangDataCollectorsEmitter.cpp
|
||||
ClangDiagnosticsEmitter.cpp
|
||||
ClangOpenCLBuiltinEmitter.cpp
|
||||
ClangOptionDocEmitter.cpp
|
||||
ClangSACheckersEmitter.cpp
|
||||
NeonEmitter.cpp
|
||||
|
||||
318
clang/utils/TableGen/ClangOpenCLBuiltinEmitter.cpp
Normal file
318
clang/utils/TableGen/ClangOpenCLBuiltinEmitter.cpp
Normal file
@ -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
|
||||
@ -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;
|
||||
|
||||
@ -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);
|
||||
|
||||
|
||||
Loading…
Reference in New Issue
Block a user