From cb18487a4449c2dce6ea71307ecdb51482d2ba84 Mon Sep 17 00:00:00 2001 From: Victor Lomuller Date: Wed, 18 Mar 2020 14:30:52 +0000 Subject: [PATCH 1/5] [SPIR-V] Make the OpenCL Builtin lookup agnostic to programming models Rename the OpenCL Builtin lookup system to make it agnostic. The Builtin emitter is now outputting the builtin information into a class so that different programming models can be produced and used. Signed-off-by: Victor Lomuller --- clang/lib/Sema/OpenCLBuiltins.td | 2 + clang/lib/Sema/SemaLookup.cpp | 127 ++++++------ clang/utils/TableGen/CMakeLists.txt | 2 +- ...r.cpp => ClangProgModelBuiltinEmitter.cpp} | 186 +++++++++++------- 4 files changed, 183 insertions(+), 134 deletions(-) rename clang/utils/TableGen/{ClangOpenCLBuiltinEmitter.cpp => ClangProgModelBuiltinEmitter.cpp} (84%) diff --git a/clang/lib/Sema/OpenCLBuiltins.td b/clang/lib/Sema/OpenCLBuiltins.td index 0d8764a567a4a..ab3310ac89204 100644 --- a/clang/lib/Sema/OpenCLBuiltins.td +++ b/clang/lib/Sema/OpenCLBuiltins.td @@ -234,6 +234,8 @@ class Builtin _Signature, list _Attributes = Attr. bit IsConst = _Attributes[1]; // Function attribute __attribute__((convergent)) bit IsConv = _Attributes[2]; + // Is function a variadic one + bit IsVariadic = 0; // OpenCL extensions to which the function belongs. FunctionExtension Extension = FuncExtNone; // Version of OpenCL from which the function is available (e.g.: CL10). diff --git a/clang/lib/Sema/SemaLookup.cpp b/clang/lib/Sema/SemaLookup.cpp index 82a197196576b..613c20c4ae76e 100644 --- a/clang/lib/Sema/SemaLookup.cpp +++ b/clang/lib/Sema/SemaLookup.cpp @@ -677,10 +677,10 @@ LLVM_DUMP_METHOD void LookupResult::dump() { D->dump(); } -/// Get the QualType instances of the return type and arguments for an OpenCL +/// Get the QualType instances of the return type and arguments for a ProgModel /// builtin function signature. /// \param Context (in) The Context instance. -/// \param OpenCLBuiltin (in) The signature currently handled. +/// \param Builtin (in) The signature currently handled. /// \param GenTypeMaxCnt (out) Maximum number of types contained in a generic /// type used as return type or as argument. /// Only meaningful for generic types, otherwise equals 1. @@ -688,27 +688,31 @@ LLVM_DUMP_METHOD void LookupResult::dump() { /// \param ArgTypes (out) List of the possible argument types. For each /// argument, ArgTypes contains QualTypes for the Cartesian product /// of (vector sizes) x (types) . -static void GetQualTypesForOpenCLBuiltin( - ASTContext &Context, const OpenCLBuiltinStruct &OpenCLBuiltin, +template +static void GetQualTypesForProgModelBuiltin( + ASTContext &Context, const typename ProgModel::BuiltinStruct &Builtin, unsigned &GenTypeMaxCnt, SmallVector &RetTypes, SmallVector, 5> &ArgTypes) { // Get the QualType instances of the return types. - unsigned Sig = SignatureTable[OpenCLBuiltin.SigTableIndex]; - OCL2Qual(Context, TypeTable[Sig], RetTypes); + unsigned Sig = ProgModel::SignatureTable[Builtin.SigTableIndex]; + ProgModel::Bultin2Qual(Context, ProgModel::TypeTable[Sig], RetTypes); GenTypeMaxCnt = RetTypes.size(); // Get the QualType instances of the arguments. // First type is the return type, skip it. - for (unsigned Index = 1; Index < OpenCLBuiltin.NumTypes; Index++) { + for (unsigned Index = 1; Index < Builtin.NumTypes; Index++) { SmallVector Ty; - OCL2Qual(Context, - TypeTable[SignatureTable[OpenCLBuiltin.SigTableIndex + Index]], Ty); + ProgModel::Bultin2Qual( + Context, + ProgModel::TypeTable[ProgModel::SignatureTable[Builtin.SigTableIndex + + Index]], + Ty); GenTypeMaxCnt = (Ty.size() > GenTypeMaxCnt) ? Ty.size() : GenTypeMaxCnt; ArgTypes.push_back(std::move(Ty)); } } -/// Create a list of the candidate function overloads for an OpenCL builtin +/// Create a list of the candidate function overloads for a ProgModel builtin /// function. /// \param Context (in) The ASTContext instance. /// \param GenTypeMaxCnt (in) Maximum number of types contained in a generic @@ -717,12 +721,12 @@ static void GetQualTypesForOpenCLBuiltin( /// \param FunctionList (out) List of FunctionTypes. /// \param RetTypes (in) List of the possible return types. /// \param ArgTypes (in) List of the possible types for the arguments. -static void GetOpenCLBuiltinFctOverloads( +static void GetProgModelBuiltinFctOverloads( ASTContext &Context, unsigned GenTypeMaxCnt, std::vector &FunctionList, SmallVector &RetTypes, - SmallVector, 5> &ArgTypes) { + SmallVector, 5> &ArgTypes, bool IsVariadic) { FunctionProtoType::ExtProtoInfo PI; - PI.Variadic = false; + PI.Variadic = IsVariadic; // Create FunctionTypes for each (gen)type. for (unsigned IGenType = 0; IGenType < GenTypeMaxCnt; IGenType++) { @@ -747,16 +751,17 @@ static void GetOpenCLBuiltinFctOverloads( /// \param S (in/out) The Sema instance. /// \param BIDecl (in) Description of the builtin. /// \param FDecl (in/out) FunctionDecl instance. -static void AddOpenCLExtensions(Sema &S, const OpenCLBuiltinStruct &BIDecl, +static void AddOpenCLExtensions(Sema &S, + const OpenCLBuiltin::BuiltinStruct &BIDecl, FunctionDecl *FDecl) { // Fetch extension associated with a function prototype. - StringRef E = FunctionExtensionTable[BIDecl.Extension]; + StringRef E = OpenCLBuiltin::FunctionExtensionTable[BIDecl.Extension]; if (E != "") S.setOpenCLExtensionForDecl(FDecl, E); } -/// When trying to resolve a function name, if isOpenCLBuiltin() returns a -/// non-null pair, then the name is referencing an OpenCL +/// When trying to resolve a function name, if ProgModel::isBuiltin() returns a +/// non-null pair, then the name is referencing a /// builtin function. Add all candidate signatures to the LookUpResult. /// /// \param S (in) The Sema instance. @@ -764,10 +769,13 @@ static void AddOpenCLExtensions(Sema &S, const OpenCLBuiltinStruct &BIDecl, /// \param II (in) The identifier being resolved. /// \param FctIndex (in) Starting index in the BuiltinTable. /// \param Len (in) The signature list has Len elements. -static void InsertOCLBuiltinDeclarationsFromTable(Sema &S, LookupResult &LR, - IdentifierInfo *II, - const unsigned FctIndex, - const unsigned Len) { +template +static void InsertBuiltinDeclarationsFromTable( + Sema &S, unsigned BuiltinSetVersion, LookupResult &LR, IdentifierInfo *II, + const unsigned FctIndex, const unsigned Len, + std::function + ProgModelFinalizer) { // The builtin function declaration uses generic types (gentype). bool HasGenType = false; @@ -776,45 +784,44 @@ static void InsertOCLBuiltinDeclarationsFromTable(Sema &S, LookupResult &LR, unsigned GenTypeMaxCnt; for (unsigned SignatureIndex = 0; SignatureIndex < Len; SignatureIndex++) { - const OpenCLBuiltinStruct &OpenCLBuiltin = - BuiltinTable[FctIndex + SignatureIndex]; + const typename ProgModel::BuiltinStruct &Builtin = + ProgModel::BuiltinTable[FctIndex + SignatureIndex]; ASTContext &Context = S.Context; // Ignore this BIF if its version does not match the language options. - unsigned OpenCLVersion = Context.getLangOpts().OpenCLVersion; - if (Context.getLangOpts().OpenCLCPlusPlus) - OpenCLVersion = 200; - if (OpenCLVersion < OpenCLBuiltin.MinVersion) - continue; - if ((OpenCLBuiltin.MaxVersion != 0) && - (OpenCLVersion >= OpenCLBuiltin.MaxVersion)) - continue; + if (BuiltinSetVersion) { + if (BuiltinSetVersion < Builtin.MinVersion) + continue; + if ((Builtin.MaxVersion != 0) && + (BuiltinSetVersion >= Builtin.MaxVersion)) + continue; + } SmallVector RetTypes; SmallVector, 5> ArgTypes; // Obtain QualType lists for the function signature. - GetQualTypesForOpenCLBuiltin(Context, OpenCLBuiltin, GenTypeMaxCnt, - RetTypes, ArgTypes); + GetQualTypesForProgModelBuiltin(Context, Builtin, GenTypeMaxCnt, + RetTypes, ArgTypes); if (GenTypeMaxCnt > 1) { HasGenType = true; } // Create function overload for each type combination. std::vector FunctionList; - GetOpenCLBuiltinFctOverloads(Context, GenTypeMaxCnt, FunctionList, RetTypes, - ArgTypes); + GetProgModelBuiltinFctOverloads(Context, GenTypeMaxCnt, FunctionList, + RetTypes, ArgTypes, Builtin.IsVariadic); SourceLocation Loc = LR.getNameLoc(); DeclContext *Parent = Context.getTranslationUnitDecl(); - FunctionDecl *NewOpenCLBuiltin; + FunctionDecl *NewBuiltin; for (unsigned Index = 0; Index < GenTypeMaxCnt; Index++) { - NewOpenCLBuiltin = FunctionDecl::Create( + NewBuiltin = FunctionDecl::Create( Context, Parent, Loc, Loc, II, FunctionList[Index], /*TInfo=*/nullptr, SC_Extern, false, FunctionList[Index]->isFunctionProtoType()); - NewOpenCLBuiltin->setImplicit(); + NewBuiltin->setImplicit(); // Create Decl objects for each parameter, adding them to the // FunctionDecl. @@ -823,29 +830,25 @@ static void InsertOCLBuiltinDeclarationsFromTable(Sema &S, LookupResult &LR, SmallVector ParmList; for (unsigned IParm = 0, e = FP->getNumParams(); IParm != e; ++IParm) { ParmVarDecl *Parm = ParmVarDecl::Create( - Context, NewOpenCLBuiltin, SourceLocation(), SourceLocation(), - nullptr, FP->getParamType(IParm), + Context, NewBuiltin, SourceLocation(), SourceLocation(), nullptr, + FP->getParamType(IParm), /*TInfo=*/nullptr, SC_None, nullptr); Parm->setScopeInfo(0, IParm); ParmList.push_back(Parm); } - NewOpenCLBuiltin->setParams(ParmList); + NewBuiltin->setParams(ParmList); } // Add function attributes. - if (OpenCLBuiltin.IsPure) - NewOpenCLBuiltin->addAttr(PureAttr::CreateImplicit(Context)); - if (OpenCLBuiltin.IsConst) - NewOpenCLBuiltin->addAttr(ConstAttr::CreateImplicit(Context)); - if (OpenCLBuiltin.IsConv) - NewOpenCLBuiltin->addAttr(ConvergentAttr::CreateImplicit(Context)); - - if (!S.getLangOpts().OpenCLCPlusPlus) - NewOpenCLBuiltin->addAttr(OverloadableAttr::CreateImplicit(Context)); - - AddOpenCLExtensions(S, OpenCLBuiltin, NewOpenCLBuiltin); - - LR.addDecl(NewOpenCLBuiltin); + if (Builtin.IsPure) + NewBuiltin->addAttr(PureAttr::CreateImplicit(Context)); + if (Builtin.IsConst) + NewBuiltin->addAttr(ConstAttr::CreateImplicit(Context)); + if (Builtin.IsConv) + NewBuiltin->addAttr(ConvergentAttr::CreateImplicit(Context)); + + ProgModelFinalizer(Builtin, *NewBuiltin); + LR.addDecl(NewBuiltin); } } @@ -878,10 +881,20 @@ bool Sema::LookupBuiltin(LookupResult &R) { // Check if this is an OpenCL Builtin, and if so, insert its overloads. if (getLangOpts().OpenCL && getLangOpts().DeclareOpenCLBuiltins) { - auto Index = isOpenCLBuiltin(II->getName()); + auto Index = OpenCLBuiltin::isBuiltin(II->getName()); if (Index.first) { - InsertOCLBuiltinDeclarationsFromTable(*this, R, II, Index.first - 1, - Index.second); + unsigned OpenCLVersion = Context.getLangOpts().OpenCLVersion; + if (Context.getLangOpts().OpenCLCPlusPlus) + OpenCLVersion = 200; + InsertBuiltinDeclarationsFromTable( + *this, OpenCLVersion, R, II, Index.first - 1, Index.second, + [this](const OpenCLBuiltin::BuiltinStruct &OpenCLBuiltin, + FunctionDecl &NewOpenCLBuiltin) { + if (!this->getLangOpts().OpenCLCPlusPlus) + NewOpenCLBuiltin.addAttr( + OverloadableAttr::CreateImplicit(Context)); + AddOpenCLExtensions(*this, OpenCLBuiltin, &NewOpenCLBuiltin); + }); return true; } } diff --git a/clang/utils/TableGen/CMakeLists.txt b/clang/utils/TableGen/CMakeLists.txt index 7deca9971090b..1d992fd04b1b2 100644 --- a/clang/utils/TableGen/CMakeLists.txt +++ b/clang/utils/TableGen/CMakeLists.txt @@ -11,7 +11,7 @@ add_tablegen(clang-tblgen CLANG ClangDataCollectorsEmitter.cpp ClangDiagnosticsEmitter.cpp ClangOpcodesEmitter.cpp - ClangOpenCLBuiltinEmitter.cpp + ClangProgModelBuiltinEmitter.cpp ClangOptionDocEmitter.cpp ClangSACheckersEmitter.cpp ClangTypeNodesEmitter.cpp diff --git a/clang/utils/TableGen/ClangOpenCLBuiltinEmitter.cpp b/clang/utils/TableGen/ClangProgModelBuiltinEmitter.cpp similarity index 84% rename from clang/utils/TableGen/ClangOpenCLBuiltinEmitter.cpp rename to clang/utils/TableGen/ClangProgModelBuiltinEmitter.cpp index b930f2daed42f..83e683d147c3f 100644 --- a/clang/utils/TableGen/ClangOpenCLBuiltinEmitter.cpp +++ b/clang/utils/TableGen/ClangProgModelBuiltinEmitter.cpp @@ -1,4 +1,4 @@ -//===- ClangOpenCLBuiltinEmitter.cpp - Generate Clang OpenCL Builtin handling +//===- ClangProgModelBuiltinEmitter.cpp - Generate Clang Builtin handling // // The LLVM Compiler Infrastructure // @@ -8,19 +8,19 @@ // //===----------------------------------------------------------------------===// // -// 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 +// This tablegen backend emits code for checking whether a function is a +// builtin function of a programming model. 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") +// For a successful lookup of e.g. the "cos" builtin, isBuiltin("cos") // returns a pair . // BuiltinTable[Index] to BuiltinTable[Index + Len] contains the pairs // of the overloads of "cos". // SignatureTable[SigIndex] to SignatureTable[SigIndex + SigLen] contains // one of the signatures of "cos". The SignatureTable entry can be // referenced by other functions, e.g. "sin", to exploit the fact that -// many OpenCL builtins share the same signature. +// many builtins may share the same signature. // // The file generated by this TableGen emitter contains the following: // @@ -31,7 +31,7 @@ // entry in this table when the builtin requires a particular (set of) // extension(s) to be enabled. // -// * OpenCLTypeStruct TypeTable[] +// * ProgModelTypeStruct TypeTable[] // Type information for return types and arguments. // // * unsigned SignatureTable[] @@ -40,17 +40,17 @@ // signature, where the first entry is the return type and subsequent // entries are the argument types. // -// * OpenCLBuiltinStruct BuiltinTable[] -// Each entry represents one overload of an OpenCL builtin function and +// * BuiltinStruct BuiltinTable[] +// Each entry represents one overload of a builtin function and // consists of an index into the SignatureTable and the number of arguments. // -// * std::pair isOpenCLBuiltin(llvm::StringRef Name) -// Find out whether a string matches an existing OpenCL builtin function +// * std::pair isBuiltin(llvm::StringRef Name) +// Find out whether a string matches an existing builtin function // name and return an index into BuiltinTable and the number of overloads. // -// * void OCL2Qual(ASTContext&, OpenCLTypeStruct, std::vector&) -// Convert an OpenCLTypeStruct type to a list of QualType instances. -// One OpenCLTypeStruct can represent multiple types, primarily when using +// * void Bultin2Qual(ASTContext&, ProgModelTypeStruct, std::vector&) +// Convert an ProgModelTypeStruct type to a list of QualType instances. +// One ProgModelTypeStruct can represent multiple types, primarily when using // GenTypes. // //===----------------------------------------------------------------------===// @@ -83,24 +83,32 @@ struct BuiltinTableEntries { class BuiltinNameEmitter { public: - BuiltinNameEmitter(RecordKeeper &Records, raw_ostream &OS) - : Records(Records), OS(OS) {} + BuiltinNameEmitter(RecordKeeper &Records, raw_ostream &OS, + llvm::StringRef Family) + : Records(Records), OS(OS), Family(Family), + ClassName((Family + "Builtin").str()) {} // Entrypoint to generate the functions and structures for checking - // whether a function is an OpenCL builtin function. + // whether a function is a builtin function. void Emit(); private: // A list of indices into the builtin function table. using BuiltinIndexListTy = SmallVector; - // Contains OpenCL builtin functions and related information, stored as + // Contains 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; + // Family for which the builtin are for. + llvm::StringRef Family; + + // Family for which the builtin are for. + std::string ClassName; + // Helper function for BuiltinNameEmitter::EmitDeclarations. Generate enum // definitions in the Output string parameter, and save their Record instances // in the List parameter. @@ -121,7 +129,7 @@ class BuiltinNameEmitter { // FctOverloadMap and TypeMap. void GetOverloads(); - // Compare two lists of signatures and check that e.g. the OpenCL version, + // Compare two lists of signatures and check that e.g. the version, // function attributes, and extension are equal for each signature. // \param Candidate (in) Entry in the SignatureListMap to check. // \param SignatureList (in) List of signatures of the considered function. @@ -134,14 +142,14 @@ class BuiltinNameEmitter { // SignatureListMap. // Some builtin functions have the same list of signatures, for example the // "sin" and "cos" functions. To save space in the BuiltinTable, the - // "isOpenCLBuiltin" function will have the same output for these two + // "isBuiltin" function will have the same output for these two // function names. void GroupBySignature(); // Emit the FunctionExtensionTable that lists all function extensions. void EmitExtensionTable(); - // Emit the TypeTable containing all types used by OpenCL builtins. + // Emit the TypeTable containing all types used by the builtins. void EmitTypeTable(); // Emit the SignatureTable. This table contains all the possible signatures. @@ -155,7 +163,7 @@ class BuiltinNameEmitter { void EmitSignatureTable(); // Emit the BuiltinTable table. This table contains all the overloads of - // each function, and is a struct OpenCLBuiltinDecl. + // each function, and is a struct BuiltinDecl. // E.g.: // // 891 convert_float2_rtn // { 58, 2, 3, 100, 0 }, @@ -163,12 +171,12 @@ class BuiltinNameEmitter { // 1 argument (+1 for the return type), stored at index 58 in // the SignatureTable. This prototype requires extension "3" in the // FunctionExtensionTable. The last two values represent the minimum (1.0) - // and maximum (0, meaning no max version) OpenCL version in which this + // and maximum (0, meaning no max version) version in which this // overload is supported. void EmitBuiltinTable(); - // Emit a StringMatcher function to check whether a function name is an - // OpenCL builtin function name. + // Emit a StringMatcher function to check whether a function name is a + // builtin function name. void EmitStringMatcher(); // Emit a function returning the clang QualType instance associated with @@ -197,14 +205,14 @@ class BuiltinNameEmitter { MapVector>> FctOverloadMap; - // Contains the map of OpenCL types to their index in the TypeTable. + // Contains the map of types to their index in the TypeTable. MapVector TypeMap; - // List of OpenCL function extensions mapping extension strings to + // List of function extensions mapping extension strings to // an index into the FunctionExtensionTable. StringMap FunctionExtensionIndex; - // List of OpenCL type names in the same order as in enum OpenCLTypeID. + // List of type names in the same order as in enum TypeID. // This list does not contain generic types. std::vector TypeList; @@ -232,7 +240,8 @@ class BuiltinNameEmitter { } // namespace void BuiltinNameEmitter::Emit() { - emitSourceFileHeader("OpenCL Builtin handling", OS); + std::string Banner = (Family + " Builtin handling").str(); + emitSourceFileHeader(Banner, OS); OS << "#include \"llvm/ADT/StringRef.h\"\n"; OS << "using namespace clang;\n\n"; @@ -263,7 +272,7 @@ void BuiltinNameEmitter::ExtractEnumTypes(std::vector &Types, for (const auto *T : Types) { if (TypesSeen.find(T->getValueAsString("Name")) == TypesSeen.end()) { - SS << " OCLT_" + T->getValueAsString("Name") << ",\n"; + SS << " TID_" + T->getValueAsString("Name") << ",\n"; // Save the type names in the same order as their enum value. Note that // the Record can be a VectorType or something else, only the name is // important. @@ -275,8 +284,11 @@ void BuiltinNameEmitter::ExtractEnumTypes(std::vector &Types, } void BuiltinNameEmitter::EmitDeclarations() { + OS << "class " << ClassName << " {\n\n" + << "public:\n\n"; + // Enum of scalar type names (float, int, ...) and generic type sets. - OS << "enum OpenCLTypeID {\n"; + OS << "enum TypeID {\n"; StringMap TypesSeen; std::string GenTypeEnums; @@ -299,17 +311,17 @@ void BuiltinNameEmitter::EmitDeclarations() { // Structure definitions. OS << R"( // Image access qualifier. -enum OpenCLAccessQual : unsigned char { - OCLAQ_None, - OCLAQ_ReadOnly, - OCLAQ_WriteOnly, - OCLAQ_ReadWrite +enum AccessQual : unsigned char { + AQ_None, + AQ_ReadOnly, + AQ_WriteOnly, + AQ_ReadWrite }; // Represents a return type or argument type. -struct OpenCLTypeStruct { +struct ProgModelTypeStruct { // A type (e.g. float, int, ...). - const OpenCLTypeID ID; + const TypeID ID; // Vector size (if applicable; 0 for scalars and generic types). const unsigned VectorWidth; // 0 if the type is not a pointer. @@ -319,14 +331,14 @@ struct OpenCLTypeStruct { // 0 if the type is not volatile. const bool IsVolatile : 1; // Access qualifier. - const OpenCLAccessQual AccessQualifier; + const AccessQual AccessQualifier; // Address space of the pointer (if applicable). const LangAS AS; }; -// One overload of an OpenCL builtin function. -struct OpenCLBuiltinStruct { - // Index of the signature in the OpenCLTypeStruct table. +// One overload of a builtin function. +struct BuiltinStruct { + // Index of the signature in the ProgModelTypeStruct table. const unsigned SigTableIndex; // Entries between index SigTableIndex and (SigTableIndex + NumTypes - 1) in // the SignatureTable represent the complete signature. The first type at @@ -338,15 +350,28 @@ struct OpenCLBuiltinStruct { const bool IsConst : 1; // Function attribute __attribute__((convergent)) const bool IsConv : 1; + // 0 if the function is not variadic. + const bool IsVariadic : 1; // OpenCL extension(s) required for this overload. const unsigned short Extension; - // First OpenCL version in which this overload was introduced (e.g. CL20). + // First version in which this overload was introduced (e.g. CL20). const unsigned short MinVersion; - // First OpenCL version in which this overload was removed (e.g. CL20). + // First version in which this overload was removed (e.g. CL20). const unsigned short MaxVersion; }; +static const char *FunctionExtensionTable[]; +static const ProgModelTypeStruct TypeTable[]; +static const unsigned short SignatureTable[]; +static const BuiltinStruct BuiltinTable[]; + +static std::pair isBuiltin(llvm::StringRef Name); +static void Bultin2Qual(ASTContext &Context, const ProgModelTypeStruct &Ty, + llvm::SmallVectorImpl &QT); + )"; + + OS << "}; // class " << ClassName << "\n"; } // Verify that the combination of GenTypes in a signature is supported. @@ -431,7 +456,7 @@ void BuiltinNameEmitter::GetOverloads() { } void BuiltinNameEmitter::EmitExtensionTable() { - OS << "static const char *FunctionExtensionTable[] = {\n"; + OS << "const char * " << ClassName << "::FunctionExtensionTable[] = {\n"; unsigned Index = 0; std::vector FuncExtensions = Records.getAllDerivedDefinitions("FunctionExtension"); @@ -448,22 +473,22 @@ void BuiltinNameEmitter::EmitExtensionTable() { } void BuiltinNameEmitter::EmitTypeTable() { - OS << "static const OpenCLTypeStruct TypeTable[] = {\n"; + OS << "const " << ClassName << "::ProgModelTypeStruct " << ClassName + << "::TypeTable[] = {\n"; for (const auto &T : TypeMap) { const char *AccessQual = StringSwitch(T.first->getValueAsString("AccessQualifier")) - .Case("RO", "OCLAQ_ReadOnly") - .Case("WO", "OCLAQ_WriteOnly") - .Case("RW", "OCLAQ_ReadWrite") - .Default("OCLAQ_None"); + .Case("RO", "AQ_ReadOnly") + .Case("WO", "AQ_WriteOnly") + .Case("RW", "AQ_ReadWrite") + .Default("AQ_None"); OS << " // " << T.second << "\n" - << " {OCLT_" << T.first->getValueAsString("Name") << ", " + << " {TID_" << T.first->getValueAsString("Name") << ", " << T.first->getValueAsInt("VecWidth") << ", " << T.first->getValueAsBit("IsPointer") << ", " << T.first->getValueAsBit("IsConst") << ", " - << T.first->getValueAsBit("IsVolatile") << ", " - << AccessQual << ", " + << T.first->getValueAsBit("IsVolatile") << ", " << AccessQual << ", " << T.first->getValueAsString("AddrSpace") << "},\n"; } OS << "};\n\n"; @@ -471,9 +496,9 @@ void BuiltinNameEmitter::EmitTypeTable() { void BuiltinNameEmitter::EmitSignatureTable() { // Store a type (e.g. int, float, int2, ...). The type is stored as an index - // of a struct OpenCLType table. Multiple entries following each other form a - // signature. - OS << "static const unsigned short SignatureTable[] = {\n"; + // of a struct ProgModelTypeStruct table. Multiple entries following each + // other form a signature. + OS << "const unsigned short " << ClassName << "::SignatureTable[] = {\n"; for (const auto &P : SignaturesList) { OS << " // " << P.second << "\n "; for (const Record *R : P.first) { @@ -494,7 +519,8 @@ void BuiltinNameEmitter::EmitSignatureTable() { void BuiltinNameEmitter::EmitBuiltinTable() { unsigned Index = 0; - OS << "static const OpenCLBuiltinStruct BuiltinTable[] = {\n"; + OS << "const " << ClassName << "::BuiltinStruct " << ClassName + << "::BuiltinTable[] = {\n"; for (const auto &SLM : SignatureListMap) { OS << " // " << (Index + 1) << ": "; @@ -510,6 +536,7 @@ void BuiltinNameEmitter::EmitBuiltinTable() { << (Overload.first->getValueAsBit("IsPure")) << ", " << (Overload.first->getValueAsBit("IsConst")) << ", " << (Overload.first->getValueAsBit("IsConv")) << ", " + << (Overload.first->getValueAsBit("IsVariadic")) << ", " << FunctionExtensionIndex[ExtName] << ", " << Overload.first->getValueAsDef("MinVersion")->getValueAsInt("ID") << ", " @@ -535,6 +562,7 @@ bool BuiltinNameEmitter::CanReuseSignature( if (Rec->getValueAsBit("IsPure") == Rec2->getValueAsBit("IsPure") && Rec->getValueAsBit("IsConst") == Rec2->getValueAsBit("IsConst") && Rec->getValueAsBit("IsConv") == Rec2->getValueAsBit("IsConv") && + Rec->getValueAsBit("IsVariadic") == Rec2->getValueAsBit("IsVariadic") && Rec->getValueAsDef("MinVersion")->getValueAsInt("ID") == Rec2->getValueAsDef("MinVersion")->getValueAsInt("ID") && Rec->getValueAsDef("MaxVersion")->getValueAsInt("ID") == @@ -611,24 +639,25 @@ void BuiltinNameEmitter::EmitStringMatcher() { } OS << R"( -// Find out whether a string matches an existing OpenCL builtin function name. +// Find out whether a string matches an existing builtin function name. // Returns: A pair <0, 0> if no name matches. // A pair indexing the BuiltinTable if the name is -// matching an OpenCL builtin function. -static std::pair isOpenCLBuiltin(llvm::StringRef Name) { - +// matching a builtin function. )"; + OS << "std::pair " << ClassName + << "::isBuiltin(llvm::StringRef Name) {\n\n"; + StringMatcher("Name", ValidBuiltins, OS).Emit(0, true); OS << " return std::make_pair(0, 0);\n"; - OS << "} // isOpenCLBuiltin\n"; + OS << "} // isBuiltin\n"; } void BuiltinNameEmitter::EmitQualTypeFinder() { OS << R"( -// Convert an OpenCLTypeStruct type to a list of QualTypes. +// Convert an ProgModelTypeStruct type to a list of QualTypes. // Generic types represent multiple types and vector sizes, thus a vector // is returned. The conversion is done in two steps: // Step 1: A switch statement fills a vector with scalar base types for the @@ -636,8 +665,13 @@ void BuiltinNameEmitter::EmitQualTypeFinder() { // or a single scalar type for non generic types. // Step 2: Qualifiers and other type properties such as vector size are // applied. -static void OCL2Qual(ASTContext &Context, const OpenCLTypeStruct &Ty, - llvm::SmallVectorImpl &QT) { +)"; + + OS << "void " << ClassName + << "::Bultin2Qual(ASTContext &Context, const ProgModelTypeStruct &Ty, " + "llvm::SmallVectorImpl &QT) {\n"; + + OS << R"( // Number of scalar types in the GenType. unsigned GenTypeNumTypes; // Pointer to the list of vector sizes for the GenType. @@ -681,16 +715,16 @@ static void OCL2Qual(ASTContext &Context, const OpenCLTypeStruct &Ty, // tells which one is needed. Emit a switch statement that puts the // corresponding QualType into "QT". for (const auto &ITE : ImageTypesMap) { - OS << " case OCLT_" << ITE.first.str() << ":\n" + OS << " case TID_" << ITE.first.str() << ":\n" << " switch (Ty.AccessQualifier) {\n" - << " case OCLAQ_None:\n" + << " case AQ_None:\n" << " llvm_unreachable(\"Image without access qualifier\");\n"; for (const auto &Image : ITE.second) { OS << StringSwitch( Image->getValueAsString("AccessQualifier")) - .Case("RO", " case OCLAQ_ReadOnly:\n") - .Case("WO", " case OCLAQ_WriteOnly:\n") - .Case("RW", " case OCLAQ_ReadWrite:\n") + .Case("RO", " case AQ_ReadOnly:\n") + .Case("WO", " case AQ_WriteOnly:\n") + .Case("RW", " case AQ_ReadWrite:\n") << " QT.push_back(Context." << Image->getValueAsDef("QTName")->getValueAsString("Name") << ");\n" << " break;\n"; @@ -701,7 +735,7 @@ static void OCL2Qual(ASTContext &Context, const OpenCLTypeStruct &Ty, // Switch cases for generic types. for (const auto *GenType : Records.getAllDerivedDefinitions("GenericType")) { - OS << " case OCLT_" << GenType->getValueAsString("Name") << ":\n"; + OS << " case TID_" << GenType->getValueAsString("Name") << ":\n"; OS << " QT.append({"; // Build the Cartesian product of (vector sizes) x (types). Only insert @@ -752,7 +786,7 @@ static void OCL2Qual(ASTContext &Context, const OpenCLTypeStruct &Ty, if (QT->getValueAsBit("IsAbstract") == 1) continue; // Emit the cases for non generic, non image types. - OS << " case OCLT_" << T->getValueAsString("Name") << ":\n"; + OS << " case TID_" << T->getValueAsString("Name") << ":\n"; OS << " QT.push_back(Context." << QT->getValueAsString("Name") << ");\n"; OS << " break;\n"; @@ -810,11 +844,11 @@ static void OCL2Qual(ASTContext &Context, const OpenCLTypeStruct &Ty, } )"; - // End of the "OCL2Qual" function. - OS << "\n} // OCL2Qual\n"; + // End of the "Bultin2Qual" function. + OS << "\n} // Bultin2Qual\n"; } void clang::EmitClangOpenCLBuiltins(RecordKeeper &Records, raw_ostream &OS) { - BuiltinNameEmitter NameChecker(Records, OS); + BuiltinNameEmitter NameChecker(Records, OS, "OpenCL"); NameChecker.Emit(); } From def61ca61c05dd593432ee94f11eae3635ec52ba Mon Sep 17 00:00:00 2001 From: Victor Lomuller Date: Wed, 18 Mar 2020 16:23:08 +0000 Subject: [PATCH 2/5] [SPIR-V] Enable SPIR-V builtin lookup Add flag -fdeclare-spirv-builtins to enable lookup of SPIR-V builtins. If -fdeclare-spirv-builtins is passed to clang, the compiler will try to lookup for the builtin described in SPIRVBuiltins.td for any match. If a match is found, overloads are build for the match. Signed-off-by: Victor Lomuller --- clang/include/clang/Basic/LangOptions.def | 1 + clang/include/clang/Driver/CC1Options.td | 2 + clang/lib/Frontend/CompilerInvocation.cpp | 1 + clang/lib/Sema/CMakeLists.txt | 6 + clang/lib/Sema/SPIRVBuiltins.td | 307 ++++++++++++++++++ clang/lib/Sema/SemaLookup.cpp | 16 + .../SemaSYCL/spirv-builtin-lookup-invalid.cpp | 12 + clang/test/SemaSYCL/spirv-builtin-lookup.cpp | 12 + .../TableGen/ClangProgModelBuiltinEmitter.cpp | 5 + clang/utils/TableGen/TableGen.cpp | 6 + clang/utils/TableGen/TableGenBackends.h | 1 + 11 files changed, 369 insertions(+) create mode 100644 clang/lib/Sema/SPIRVBuiltins.td create mode 100644 clang/test/SemaSYCL/spirv-builtin-lookup-invalid.cpp create mode 100644 clang/test/SemaSYCL/spirv-builtin-lookup.cpp diff --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def index 631dc435b97f8..9ff8283702970 100644 --- a/clang/include/clang/Basic/LangOptions.def +++ b/clang/include/clang/Basic/LangOptions.def @@ -237,6 +237,7 @@ LANGOPT(SYCLAllowFuncPtr , 1, 0, "Allow function pointers in SYCL device code") LANGOPT(SYCLStdLayoutKernelParams, 1, 0, "Enable standard layout requirement for SYCL kernel parameters") LANGOPT(SYCLUnnamedLambda , 1, 0, "Allow unnamed lambda SYCL kernels") LANGOPT(SYCLVersion , 32, 0, "Version of the SYCL standard used") +LANGOPT(DeclareSPIRVBuiltins, 1, 0, "Declare SPIR-V builtin functions") LANGOPT(HIPUseNewLaunchAPI, 1, 0, "Use new kernel launching API for HIP") diff --git a/clang/include/clang/Driver/CC1Options.td b/clang/include/clang/Driver/CC1Options.td index 3f223609f3f8d..8998fc6271e7b 100644 --- a/clang/include/clang/Driver/CC1Options.td +++ b/clang/include/clang/Driver/CC1Options.td @@ -807,6 +807,8 @@ def finclude_default_header : Flag<["-"], "finclude-default-header">, HelpText<"Include default header file for OpenCL">; def fdeclare_opencl_builtins : Flag<["-"], "fdeclare-opencl-builtins">, HelpText<"Add OpenCL builtin function declarations (experimental)">; +def fdeclare_spirv_builtins : Flag<["-"], "fdeclare-spirv-builtins">, + HelpText<"Add SPIR-V 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 24053edff5071..fdfd5a0fafc24 100644 --- a/clang/lib/Frontend/CompilerInvocation.cpp +++ b/clang/lib/Frontend/CompilerInvocation.cpp @@ -2577,6 +2577,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); + Opts.DeclareSPIRVBuiltins = Args.hasArg(OPT_fdeclare_spirv_builtins); llvm::Triple T(TargetOpts.Triple); CompilerInvocation::setLangDefaults(Opts, IK, T, PPOpts, LangStd); diff --git a/clang/lib/Sema/CMakeLists.txt b/clang/lib/Sema/CMakeLists.txt index 75255f990d80e..f3e5f3cc632e6 100644 --- a/clang/lib/Sema/CMakeLists.txt +++ b/clang/lib/Sema/CMakeLists.txt @@ -17,6 +17,11 @@ clang_tablegen(OpenCLBuiltins.inc -gen-clang-opencl-builtins TARGET ClangOpenCLBuiltinsImpl ) +clang_tablegen(SPIRVBuiltins.inc -gen-clang-spirv-builtins + SOURCE SPIRVBuiltins.td + TARGET ClangSPIRVBuiltinsImpl + ) + add_clang_library(clangSema AnalysisBasedWarnings.cpp CodeCompleteConsumer.cpp @@ -72,6 +77,7 @@ add_clang_library(clangSema DEPENDS ClangOpenCLBuiltinsImpl + ClangSPIRVBuiltinsImpl LINK_LIBS clangAST diff --git a/clang/lib/Sema/SPIRVBuiltins.td b/clang/lib/Sema/SPIRVBuiltins.td new file mode 100644 index 0000000000000..f11f97685e0f5 --- /dev/null +++ b/clang/lib/Sema/SPIRVBuiltins.td @@ -0,0 +1,307 @@ +//==--- SPIRVBuiltins.td - SPIRV 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 SPIR-V builtin function +// declarations. In case of an unresolved function name, Clang will check for +// a function described in this file when -fdeclare-spirv-builtins is specified. +// +//===----------------------------------------------------------------------===// + +//===----------------------------------------------------------------------===// +// Definitions of miscellaneous basic entities. +//===----------------------------------------------------------------------===// +// TODO: basic entities declaration with OpenCLBuiltins.td + +// TODO: Manage version using the JSON grammar. Unused for now. +class Version { + int ID = _Version; +} +def SPIRVAll : Version< 0>; + +// Address spaces +// Pointer types need to be assigned an address space. +class AddressSpace { + string Name = _AS; +} +def DefaultAS : AddressSpace<"clang::LangAS::Default">; +def PrivateAS : AddressSpace<"clang::LangAS::sycl_private">; +def GlobalAS : AddressSpace<"clang::LangAS::sycl_global">; +def ConstantAS : AddressSpace<"clang::LangAS::sycl_constant">; +def LocalAS : AddressSpace<"clang::LangAS::sycl_local">; +def GenericAS : AddressSpace<"clang::LangAS::sycl_generic">; + +// TODO: Manage capabilities. Unused for now. +class AbstractExtension { + string ExtName = _Ext; +} + +// Extension associated to a builtin function. +class FunctionExtension : AbstractExtension<_Ext>; + +// FunctionExtension definitions. +def FuncExtNone : FunctionExtension<"">; + +// Qualified Type. These map to 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; + // Some QualTypes in this file represent an abstract type for which there is + // no corresponding AST QualType, e.g. a GenType or an `image2d_t` type + // without access qualifiers. + bit IsAbstract = _IsAbstract; + bit IsSigned = _IsSigned; +} + +// List of integers. +class IntList _List> { + string Name = _Name; + list List = _List; +} + +// Basic data types (int, float, image2d_t, ...). +// Its child classes can represent concrete types (e.g. VectorType) or +// abstract types (e.g. GenType). +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 = 1; + // Size of the element in bits. + int ElementSize = 1; + // Is a integer. + bit IsInteger = 0; + // Is a signed integer. + bit IsSigned = 1; + // Is a float. + bit IsFloat = 0; + // Is a pointer. + bit IsPointer = 0; + // "const" qualifier. + bit IsConst = 0; + // "volatile" qualifier. + bit IsVolatile = 0; + // Access qualifier. Must be one of ("RO", "WO", "RW"). + string AccessQualifier = ""; + // Address space. + string AddrSpace = DefaultAS.Name; +} + +class FundamentalType : Type<_Name, _QTName> { + // Inherited fields + let ElementSize = _Size; +} + +// Integer Type. +class IntType : FundamentalType<_Name, _QTName, _Size> { + // Inherited fields + let IsInteger = 1; + let IsSigned = 1; +} + +// Unsigned integer Type. +class UIntType : FundamentalType<_Name, _QTName, _Size> { + // Inherited fields + let IsInteger = 1; + let IsSigned = 0; +} + +// Floating Type. +class FPType : FundamentalType<_Name, _QTName, _Size> { + // Inherited fields + let IsFloat = 1; +} + +class CompoundType : Type<_Ty.Name, _Ty.QTName> { + // Inherited fields + let VecWidth = _Ty.VecWidth; + let ElementSize = _Ty.ElementSize; + let IsInteger = _Ty.IsInteger; + let IsSigned = _Ty.IsSigned; + let IsFloat = _Ty.IsFloat; + let IsPointer = _Ty.IsPointer; + let IsConst = _Ty.IsConst; + let IsVolatile = _Ty.IsVolatile; + let AccessQualifier = _Ty.AccessQualifier; + let AddrSpace = _Ty.AddrSpace; + + Type ElementType = _Ty; +} + +// Vector types (e.g. int2, int3, int16, float8, ...). +class VectorType : Type<_Ty.Name, _Ty.QTName> { + let VecWidth = _VecWidth; + let AccessQualifier = ""; + // Inherited fields + let ElementSize = _Ty.ElementSize; + let IsInteger = _Ty.IsInteger; + let IsSigned = _Ty.IsSigned; + let IsFloat = _Ty.IsFloat; + let IsPointer = _Ty.IsPointer; + let IsConst = _Ty.IsConst; + let IsVolatile = _Ty.IsVolatile; + let AccessQualifier = _Ty.AccessQualifier; + let AddrSpace = _Ty.AddrSpace; +} + +// Pointer types (e.g. int*, float*, ...). +class PointerType : + CompoundType<_Ty> { + // Inherited fields + let IsPointer = 1; + let AddrSpace = _AS.Name; +} + +// Const types (e.g. const int). +class ConstType : CompoundType<_Ty> { + // Inherited fields + let IsConst = 1; +} + +// Volatile types (e.g. volatile int). +class VolatileType : CompoundType<_Ty> { + // Inherited fields + let IsVolatile = 1; +} + +// Image types (e.g. image2d). +class ImageType : + Type<_Ty.Name, QualType<_Ty.QTName.Name#_AccessQualifier#"Ty", 0>> { + let VecWidth = 0; + let AccessQualifier = _AccessQualifier; + // Inherited fields + let ElementSize = _Ty.ElementSize; + let IsInteger = _Ty.IsInteger; + let IsSigned = _Ty.IsSigned; + let IsFloat = _Ty.IsFloat; + let IsPointer = _Ty.IsPointer; + let IsConst = _Ty.IsConst; + let IsVolatile = _Ty.IsVolatile; + let AddrSpace = _Ty.AddrSpace; +} + +// List of Types. +class TypeList _Type> { + list List = _Type; +} + +// A GenericType is an abstract type that defines a set of types as a +// combination of Types and vector sizes. +// +// For example, if TypeList = and VectorList = <1, 2, 4>, then it +// represents . +// +// Some rules apply when using multiple GenericType arguments in a declaration: +// 1. The number of vector sizes must be equal or 1 for all gentypes in a +// declaration. +// 2. The number of Types must be equal or 1 for all gentypes in a +// declaration. +// 3. Generic types are combined by iterating over all generic types at once. +// For example, for the following GenericTypes +// GenT1 = GenericType and +// GenT2 = GenericType +// A declaration f(GenT1, GenT2) results in the combinations +// f(half, float), f(half2, float2), f(half, int), f(half2, int2) . +// 4. "sgentype" from the OpenCL specification is supported by specifying +// a single vector size. +// For example, for the following GenericTypes +// GenT = GenericType and +// SGenT = GenericType +// A declaration f(GenT, SGenT) results in the combinations +// f(half, half), f(half2, half), f(int, int), f(int2, int) . +class GenericType : + Type<_Ty, QualType<"null", 1>> { + // Possible element types of the generic type. + TypeList TypeList = _TypeList; + // Possible vector sizes of the types in the TypeList. + IntList VectorList = _VectorList; + // The VecWidth field is ignored for GenericTypes. Use VectorList instead. + let VecWidth = 0; +} + +// Builtin function attributes. +def Attr { + list None = [0, 0, 0]; + list Pure = [1, 0, 0]; + list Const = [0, 1, 0]; + list Convergent = [0, 0, 1]; +} + +//===----------------------------------------------------------------------===// +// Class for builtin functions +//===----------------------------------------------------------------------===// +class Builtin _Signature, list _Attributes = Attr.None> { + // 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; + // Function attribute __attribute__((pure)) + bit IsPure = _Attributes[0]; + // Function attribute __attribute__((const)) + bit IsConst = _Attributes[1]; + // Function attribute __attribute__((convergent)) + bit IsConv = _Attributes[2]; + // Is function a variadic one + bit IsVariadic = 0; + // OpenCL extensions to which the function belongs. + FunctionExtension Extension = FuncExtNone; + // Version from which the function is available. + // MinVersion is inclusive. + Version MinVersion = SPIRVAll; + // Version from which the function is not supported anymore. + // MaxVersion is exclusive. + // SPIRVAll makes the function available for all versions. + Version MaxVersion = SPIRVAll; +} + +// Helper to declare SPIR-V Core builtins. +class SPVBuiltin _Signature, list _Attributes = Attr.None> : +Builtin<"__spirv_" # _Name, _Signature, _Attributes> {} + +// Helper to declare OpenCL SPIR-V extended set builtins. +class OCLSPVBuiltin _Signature, list _Attributes = Attr.None> : +SPVBuiltin<"ocl_" # _Name, _Signature, _Attributes> {} + +//===----------------------------------------------------------------------===// +// Definitions of types +//===----------------------------------------------------------------------===// + +def Float : FPType<"float", QualType<"FloatTy">, 32>; +def Double : FPType<"double", QualType<"DoubleTy">, 64>; +def Half : FPType<"half", QualType<"Float16Ty">, 16>; + +//===----------------------------------------------------------------------===// +// Definitions of gentype variants +//===----------------------------------------------------------------------===// + +// Vector width lists. +def VecAndScalar: IntList<"VecAndScalar", [1, 2, 3, 4, 8, 16]>; + +// Type lists. +def TLFloat : TypeList<[Float, Double, Half]>; + +// Float +def FGenTypeN : GenericType<"FGenTypeN", TLFloat, VecAndScalar>; + + + +//===----------------------------------------------------------------------===// +// Definitions of builtins +// extinst.opencl.std.100.grammar.json +//===----------------------------------------------------------------------===// + +// 2.1. Math extended instructions + +def : OCLSPVBuiltin<"acos", [FGenTypeN, FGenTypeN], Attr.Const>; + diff --git a/clang/lib/Sema/SemaLookup.cpp b/clang/lib/Sema/SemaLookup.cpp index 613c20c4ae76e..6ff9d8409f3fe 100644 --- a/clang/lib/Sema/SemaLookup.cpp +++ b/clang/lib/Sema/SemaLookup.cpp @@ -48,6 +48,7 @@ #include #include "OpenCLBuiltins.inc" +#include "SPIRVBuiltins.inc" using namespace clang; using namespace sema; @@ -899,6 +900,21 @@ bool Sema::LookupBuiltin(LookupResult &R) { } } + // Check if this is a SPIR-V Builtin, and if so, insert its overloads. + if (getLangOpts().DeclareSPIRVBuiltins) { + auto Index = SPIRVBuiltin::isBuiltin(II->getName()); + if (Index.first) { + InsertBuiltinDeclarationsFromTable( + *this, 0, R, II, Index.first - 1, Index.second, + [this](const SPIRVBuiltin::BuiltinStruct &, + FunctionDecl &NewBuiltin) { + NewBuiltin.addAttr( + SYCLDeviceAttr::CreateImplicit(this->Context)); + }); + 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/SemaSYCL/spirv-builtin-lookup-invalid.cpp b/clang/test/SemaSYCL/spirv-builtin-lookup-invalid.cpp new file mode 100644 index 0000000000000..757d10dcfc8b5 --- /dev/null +++ b/clang/test/SemaSYCL/spirv-builtin-lookup-invalid.cpp @@ -0,0 +1,12 @@ +// RUN: %clang_cc1 -fdeclare-spirv-builtins -fsyntax-only -verify -std=c++11 %s + +// Verify that invalid call to __spirv_ocl_acos (no viable overloads) get diagnosed + +struct InvalidType {}; +void acos(InvalidType Invalid) { + __spirv_ocl_acos(Invalid); // expected-error {{no matching function for call to '__spirv_ocl_acos'}} + // expected-note@-1 + {{candidate function not viable: no known conversion from}} + // too many params + __spirv_ocl_acos(42.f, 42.f); // expected-error {{no matching function for call to '__spirv_ocl_acos'}} + // expected-note@-1 + {{candidate function not viable: requires 1 argument, but 2 were provided}} +} diff --git a/clang/test/SemaSYCL/spirv-builtin-lookup.cpp b/clang/test/SemaSYCL/spirv-builtin-lookup.cpp new file mode 100644 index 0000000000000..d6c8bc69a2f2f --- /dev/null +++ b/clang/test/SemaSYCL/spirv-builtin-lookup.cpp @@ -0,0 +1,12 @@ +// RUN: %clang_cc1 -fdeclare-spirv-builtins -fsyntax-only -verify -std=c++11 %s +// expected-no-diagnostics + +// Verify that __spirv_ocl_acos is recognized as a builtin + +float acos(float val) { + return __spirv_ocl_acos(val); +} + +double acos(double val) { + return __spirv_ocl_acos(val); +} diff --git a/clang/utils/TableGen/ClangProgModelBuiltinEmitter.cpp b/clang/utils/TableGen/ClangProgModelBuiltinEmitter.cpp index 83e683d147c3f..b431ce61ff689 100644 --- a/clang/utils/TableGen/ClangProgModelBuiltinEmitter.cpp +++ b/clang/utils/TableGen/ClangProgModelBuiltinEmitter.cpp @@ -852,3 +852,8 @@ void clang::EmitClangOpenCLBuiltins(RecordKeeper &Records, raw_ostream &OS) { BuiltinNameEmitter NameChecker(Records, OS, "OpenCL"); NameChecker.Emit(); } + +void clang::EmitClangSPIRVBuiltins(RecordKeeper &Records, raw_ostream &OS) { + BuiltinNameEmitter NameChecker(Records, OS, "SPIRV"); + NameChecker.Emit(); +} diff --git a/clang/utils/TableGen/TableGen.cpp b/clang/utils/TableGen/TableGen.cpp index 3d8f6dc352d01..237afb32e57ea 100644 --- a/clang/utils/TableGen/TableGen.cpp +++ b/clang/utils/TableGen/TableGen.cpp @@ -61,6 +61,7 @@ enum ActionType { GenClangCommentCommandInfo, GenClangCommentCommandList, GenClangOpenCLBuiltins, + GenClangSPIRVBuiltins, GenArmNeon, GenArmFP16, GenArmNeonSema, @@ -179,6 +180,8 @@ cl::opt Action( "documentation comments"), clEnumValN(GenClangOpenCLBuiltins, "gen-clang-opencl-builtins", "Generate OpenCL builtin declaration handlers"), + clEnumValN(GenClangSPIRVBuiltins, "gen-clang-spirv-builtins", + "Generate SPIR-V 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", @@ -339,6 +342,9 @@ bool ClangTableGenMain(raw_ostream &OS, RecordKeeper &Records) { case GenClangOpenCLBuiltins: EmitClangOpenCLBuiltins(Records, OS); break; + case GenClangSPIRVBuiltins: + EmitClangSPIRVBuiltins(Records, OS); + break; case GenArmNeon: EmitNeon(Records, OS); break; diff --git a/clang/utils/TableGen/TableGenBackends.h b/clang/utils/TableGen/TableGenBackends.h index cc300319c1006..1785180689cb1 100644 --- a/clang/utils/TableGen/TableGenBackends.h +++ b/clang/utils/TableGen/TableGenBackends.h @@ -109,6 +109,7 @@ void EmitClangOptDocs(llvm::RecordKeeper &Records, llvm::raw_ostream &OS); void EmitClangOpenCLBuiltins(llvm::RecordKeeper &Records, llvm::raw_ostream &OS); +void EmitClangSPIRVBuiltins(llvm::RecordKeeper &Records, llvm::raw_ostream &OS); void EmitClangDataCollectors(llvm::RecordKeeper &Records, llvm::raw_ostream &OS); From eb1d13232ebca50a4ffde49499ce2675a5cfad10 Mon Sep 17 00:00:00 2001 From: Victor Lomuller Date: Fri, 20 Mar 2020 12:35:33 +0000 Subject: [PATCH 3/5] Apply suggestions from code review Co-Authored-By: Alexey Bader Signed-off-by: Victor Lomuller --- clang/test/SemaSYCL/spirv-builtin-lookup-invalid.cpp | 2 +- clang/test/SemaSYCL/spirv-builtin-lookup.cpp | 2 +- clang/utils/TableGen/ClangProgModelBuiltinEmitter.cpp | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/test/SemaSYCL/spirv-builtin-lookup-invalid.cpp b/clang/test/SemaSYCL/spirv-builtin-lookup-invalid.cpp index 757d10dcfc8b5..b30dfc9ae3b6f 100644 --- a/clang/test/SemaSYCL/spirv-builtin-lookup-invalid.cpp +++ b/clang/test/SemaSYCL/spirv-builtin-lookup-invalid.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fdeclare-spirv-builtins -fsyntax-only -verify -std=c++11 %s +// RUN: %clang_cc1 -fdeclare-spirv-builtins -fsyntax-only -verify %s // Verify that invalid call to __spirv_ocl_acos (no viable overloads) get diagnosed diff --git a/clang/test/SemaSYCL/spirv-builtin-lookup.cpp b/clang/test/SemaSYCL/spirv-builtin-lookup.cpp index d6c8bc69a2f2f..45ed5c4fd8384 100644 --- a/clang/test/SemaSYCL/spirv-builtin-lookup.cpp +++ b/clang/test/SemaSYCL/spirv-builtin-lookup.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fdeclare-spirv-builtins -fsyntax-only -verify -std=c++11 %s +// RUN: %clang_cc1 -fdeclare-spirv-builtins -fsyntax-only -verify %s // expected-no-diagnostics // Verify that __spirv_ocl_acos is recognized as a builtin diff --git a/clang/utils/TableGen/ClangProgModelBuiltinEmitter.cpp b/clang/utils/TableGen/ClangProgModelBuiltinEmitter.cpp index b431ce61ff689..d7e4304327e49 100644 --- a/clang/utils/TableGen/ClangProgModelBuiltinEmitter.cpp +++ b/clang/utils/TableGen/ClangProgModelBuiltinEmitter.cpp @@ -106,7 +106,7 @@ class BuiltinNameEmitter { // Family for which the builtin are for. llvm::StringRef Family; - // Family for which the builtin are for. + // Class for which the builtin are for. std::string ClassName; // Helper function for BuiltinNameEmitter::EmitDeclarations. Generate enum From 270933999abc1a3839b551ce579687a5c09c1b5a Mon Sep 17 00:00:00 2001 From: Victor Lomuller Date: Fri, 20 Mar 2020 12:55:14 +0000 Subject: [PATCH 4/5] [SPIR-V] Add CodeGen test for the SPIR-V builtin lookup Signed-off-by: Victor Lomuller --- .../test/CodeGenSPIRV/spirv-builtin-lookup.cpp | 17 +++++++++++++++++ 1 file changed, 17 insertions(+) create mode 100644 clang/test/CodeGenSPIRV/spirv-builtin-lookup.cpp diff --git a/clang/test/CodeGenSPIRV/spirv-builtin-lookup.cpp b/clang/test/CodeGenSPIRV/spirv-builtin-lookup.cpp new file mode 100644 index 0000000000000..438a226f39604 --- /dev/null +++ b/clang/test/CodeGenSPIRV/spirv-builtin-lookup.cpp @@ -0,0 +1,17 @@ +// RUN: %clang_cc1 -fdeclare-spirv-builtins -fsyntax-only -emit-llvm %s -o - | FileCheck %s + +float acos(float val) { + // CHECK: @_Z4acosf + // CHECK: call float @_Z16__spirv_ocl_acosf + return __spirv_ocl_acos(val); +} + +// CHECK: declare float @_Z16__spirv_ocl_acosf(float) + +double acos(double val) { + // CHECK: @_Z4acosd + // CHECK: call double @_Z16__spirv_ocl_acosd + return __spirv_ocl_acos(val); +} + +// CHECK: declare double @_Z16__spirv_ocl_acosd(double) From 289fffe8c58ce933b437de0edf68e9a1e1536aee Mon Sep 17 00:00:00 2001 From: Victor Lomuller Date: Fri, 20 Mar 2020 18:49:18 +0000 Subject: [PATCH 5/5] [SPIR-V] Add test with windows mangling Signed-off-by: Victor Lomuller --- .../CodeGenSPIRV/spirv-builtin-lookup-win.cpp | 17 +++++++++++++++++ .../test/CodeGenSPIRV/spirv-builtin-lookup.cpp | 2 +- 2 files changed, 18 insertions(+), 1 deletion(-) create mode 100644 clang/test/CodeGenSPIRV/spirv-builtin-lookup-win.cpp diff --git a/clang/test/CodeGenSPIRV/spirv-builtin-lookup-win.cpp b/clang/test/CodeGenSPIRV/spirv-builtin-lookup-win.cpp new file mode 100644 index 0000000000000..1d477a20dc6fc --- /dev/null +++ b/clang/test/CodeGenSPIRV/spirv-builtin-lookup-win.cpp @@ -0,0 +1,17 @@ +// RUN: %clang_cc1 -triple x86_64-windows-msvc -fdeclare-spirv-builtins -fsyntax-only -emit-llvm %s -o - | FileCheck %s + +float acos(float val) { + // CHECK: @"?acos@@YAMM@Z" + // CHECK: call float @"?__spirv_ocl_acos@@YAMM@Z" + return __spirv_ocl_acos(val); +} + +// CHECK: declare dso_local float @"?__spirv_ocl_acos@@YAMM@Z"(float) + +double acos(double val) { + // CHECK: @"?acos@@YANN@Z" + // CHECK: call double @"?__spirv_ocl_acos@@YANN@Z" + return __spirv_ocl_acos(val); +} + +// CHECK: declare dso_local double @"?__spirv_ocl_acos@@YANN@Z"(double) diff --git a/clang/test/CodeGenSPIRV/spirv-builtin-lookup.cpp b/clang/test/CodeGenSPIRV/spirv-builtin-lookup.cpp index 438a226f39604..a6805c12aa55e 100644 --- a/clang/test/CodeGenSPIRV/spirv-builtin-lookup.cpp +++ b/clang/test/CodeGenSPIRV/spirv-builtin-lookup.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fdeclare-spirv-builtins -fsyntax-only -emit-llvm %s -o - | FileCheck %s +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fdeclare-spirv-builtins -fsyntax-only -emit-llvm %s -o - | FileCheck %s float acos(float val) { // CHECK: @_Z4acosf