Index: clang/utils/TableGen/ClangOpenCLBuiltinEmitter.cpp =================================================================== --- clang/utils/TableGen/ClangOpenCLBuiltinEmitter.cpp +++ clang/utils/TableGen/ClangOpenCLBuiltinEmitter.cpp @@ -229,19 +229,17 @@ MapVector SignatureListMap; }; -// OpenCL builtin test generator. This class processes the same TableGen input -// as BuiltinNameEmitter, but generates a .cl file that contains a call to each -// builtin function described in the .td input. -class OpenCLBuiltinTestEmitter { +/// Base class for emitting a file (e.g. header or test) from OpenCLBuiltins.td +class OpenCLBuiltinFileEmitterBase { public: - OpenCLBuiltinTestEmitter(RecordKeeper &Records, raw_ostream &OS) + OpenCLBuiltinFileEmitterBase(RecordKeeper &Records, raw_ostream &OS) : Records(Records), OS(OS) {} // Entrypoint to generate the functions for testing all OpenCL builtin // functions. - void emit(); + virtual void emit() = 0; -private: +protected: struct TypeFlags { TypeFlags() : IsConst(false), IsVolatile(false), IsPointer(false) {} bool IsConst : 1; @@ -278,6 +276,18 @@ expandTypesInSignature(const std::vector &Signature, SmallVectorImpl> &Types); + // Emit extension enabling pragmas. + void emitExtensionSetup(); + + // Emit an #if guard for a Builtin's extension. Return the corresponding + // closing #endif, or an empty string if no extension #if guard was emitted. + std::string emitExtensionGuard(const Record *Builtin); + + // Emit an #if guard for a Builtin's language version. Return the + // corresponding closing #endif, or an empty string if no version #if guard + // was emitted. + std::string emitVersionGuard(const Record *Builtin); + // Contains OpenCL builtin functions and related information, stored as // Record instances. They are coming from the associated TableGen file. RecordKeeper &Records; @@ -286,6 +296,31 @@ raw_ostream &OS; }; +// OpenCL builtin test generator. This class processes the same TableGen input +// as BuiltinNameEmitter, but generates a .cl file that contains a call to each +// builtin function described in the .td input. +class OpenCLBuiltinTestEmitter : public OpenCLBuiltinFileEmitterBase { +public: + OpenCLBuiltinTestEmitter(RecordKeeper &Records, raw_ostream &OS) + : OpenCLBuiltinFileEmitterBase(Records, OS) {} + + // Entrypoint to generate the functions for testing all OpenCL builtin + // functions. + void emit() override; +}; + +// OpenCL builtin header generator. This class processes the same TableGen +// input as BuiltinNameEmitter, but generates a .h file that contains a +// prototype for each builtin function described in the .td input. +class OpenCLBuiltinHeaderEmitter : public OpenCLBuiltinFileEmitterBase { +public: + OpenCLBuiltinHeaderEmitter(RecordKeeper &Records, raw_ostream &OS) + : OpenCLBuiltinFileEmitterBase(Records, OS) {} + + // Entrypoint to generate the header. + void emit() override; +}; + } // namespace void BuiltinNameEmitter::Emit() { @@ -919,9 +954,9 @@ OS << "\n} // OCL2Qual\n"; } -std::string OpenCLBuiltinTestEmitter::getTypeString(const Record *Type, - TypeFlags Flags, - int VectorSize) const { +std::string OpenCLBuiltinFileEmitterBase::getTypeString(const Record *Type, + TypeFlags Flags, + int VectorSize) const { std::string S; if (Type->getValueAsBit("IsConst") || Flags.IsConst) { S += "const "; @@ -966,7 +1001,7 @@ return S; } -void OpenCLBuiltinTestEmitter::getTypeLists( +void OpenCLBuiltinFileEmitterBase::getTypeLists( Record *Type, TypeFlags &Flags, std::vector &TypeList, std::vector &VectorList) const { bool isGenType = Type->isSubClassOf("GenericType"); @@ -999,7 +1034,7 @@ VectorList.push_back(Type->getValueAsInt("VecWidth")); } -void OpenCLBuiltinTestEmitter::expandTypesInSignature( +void OpenCLBuiltinFileEmitterBase::expandTypesInSignature( const std::vector &Signature, SmallVectorImpl> &Types) { // Find out if there are any GenTypes in this signature, and if so, calculate @@ -1040,10 +1075,7 @@ } } -void OpenCLBuiltinTestEmitter::emit() { - emitSourceFileHeader("OpenCL Builtin exhaustive testing", OS); - - // Enable some extensions for testing. +void OpenCLBuiltinFileEmitterBase::emitExtensionSetup() { OS << R"( #pragma OPENCL EXTENSION cl_khr_fp16 : enable #pragma OPENCL EXTENSION cl_khr_fp64 : enable @@ -1054,6 +1086,60 @@ #pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable )"; +} + +std::string +OpenCLBuiltinFileEmitterBase::emitExtensionGuard(const Record *Builtin) { + StringRef Extensions = + Builtin->getValueAsDef("Extension")->getValueAsString("ExtName"); + if (Extensions.empty()) + return ""; + + OS << "#if"; + + SmallVector ExtVec; + Extensions.split(ExtVec, " "); + bool isFirst = true; + for (StringRef Ext : ExtVec) { + if (!isFirst) { + OS << " &&"; + } + OS << " defined(" << Ext << ")"; + isFirst = false; + } + OS << "\n"; + + return "#endif // Extension\n"; +} + +std::string +OpenCLBuiltinFileEmitterBase::emitVersionGuard(const Record *Builtin) { + std::string OptionalEndif; + auto PrintOpenCLVersion = [this](int Version) { + OS << "CL_VERSION_" << (Version / 100) << "_" << ((Version % 100) / 10); + }; + int MinVersion = Builtin->getValueAsDef("MinVersion")->getValueAsInt("ID"); + if (MinVersion != 100) { + // OpenCL 1.0 is the default minimum version. + OS << "#if __OPENCL_C_VERSION__ >= "; + PrintOpenCLVersion(MinVersion); + OS << "\n"; + OptionalEndif = "#endif // MinVersion\n" + OptionalEndif; + } + int MaxVersion = Builtin->getValueAsDef("MaxVersion")->getValueAsInt("ID"); + if (MaxVersion) { + OS << "#if __OPENCL_C_VERSION__ < "; + PrintOpenCLVersion(MaxVersion); + OS << "\n"; + OptionalEndif = "#endif // MaxVersion\n" + OptionalEndif; + } + return OptionalEndif; +} + +void OpenCLBuiltinTestEmitter::emit() { + emitSourceFileHeader("OpenCL Builtin exhaustive testing", OS); + + emitExtensionSetup(); // Ensure each test has a unique name by numbering them. unsigned TestID = 0; @@ -1067,43 +1153,10 @@ expandTypesInSignature(B->getValueAsListOfDefs("Signature"), FTypes); OS << "// Test " << Name << "\n"; - std::string OptionalEndif; - StringRef Extensions = - B->getValueAsDef("Extension")->getValueAsString("ExtName"); - if (!Extensions.empty()) { - OS << "#if"; - OptionalEndif = "#endif // Extension\n"; - - SmallVector ExtVec; - Extensions.split(ExtVec, " "); - bool isFirst = true; - for (StringRef Ext : ExtVec) { - if (!isFirst) { - OS << " &&"; - } - OS << " defined(" << Ext << ")"; - isFirst = false; - } - OS << "\n"; - } - auto PrintOpenCLVersion = [this](int Version) { - OS << "CL_VERSION_" << (Version / 100) << "_" << ((Version % 100) / 10); - }; - int MinVersion = B->getValueAsDef("MinVersion")->getValueAsInt("ID"); - if (MinVersion != 100) { - // OpenCL 1.0 is the default minimum version. - OS << "#if __OPENCL_C_VERSION__ >= "; - PrintOpenCLVersion(MinVersion); - OS << "\n"; - OptionalEndif = "#endif // MinVersion\n" + OptionalEndif; - } - int MaxVersion = B->getValueAsDef("MaxVersion")->getValueAsInt("ID"); - if (MaxVersion) { - OS << "#if __OPENCL_C_VERSION__ < "; - PrintOpenCLVersion(MaxVersion); - OS << "\n"; - OptionalEndif = "#endif // MaxVersion\n" + OptionalEndif; - } + + std::string OptionalExtensionEndif = emitExtensionGuard(B); + std::string OptionalVersionEndif = emitVersionGuard(B); + for (const auto &Signature : FTypes) { // Emit function declaration. OS << Signature[0] << " test" << TestID++ << "_" << Name << "("; @@ -1132,15 +1185,76 @@ // End of function body. OS << "}\n"; } - OS << OptionalEndif << "\n"; + + OS << OptionalVersionEndif; + OS << OptionalExtensionEndif; } } +void OpenCLBuiltinHeaderEmitter::emit() { + emitSourceFileHeader("OpenCL Builtin declarations", OS); + + emitExtensionSetup(); + + OS << R"( +#define __ovld __attribute__((overloadable)) +#define __conv __attribute__((convergent)) +#define __purefn __attribute__((pure)) +#define __cnfn __attribute__((const)) + +)"; + + // Iterate over all builtins. + std::vector Builtins = Records.getAllDerivedDefinitions("Builtin"); + for (const auto *B : Builtins) { + StringRef Name = B->getValueAsString("Name"); + + std::string OptionalExtensionEndif = emitExtensionGuard(B); + std::string OptionalVersionEndif = emitVersionGuard(B); + + SmallVector, 4> FTypes; + expandTypesInSignature(B->getValueAsListOfDefs("Signature"), FTypes); + + for (const auto &Signature : FTypes) { + // Emit function declaration. + OS << "__ovld "; + if (B->getValueAsBit("IsConst")) + OS << "__cnfn "; + if (B->getValueAsBit("IsPure")) + OS << "__purefn "; + if (B->getValueAsBit("IsConv")) + OS << "__conv "; + + OS << Signature[0] << " " << Name << "("; + if (Signature.size() > 1) { + for (unsigned I = 1; I < Signature.size(); I++) { + if (I != 1) + OS << ", "; + OS << Signature[I]; + } + } + OS << ");\n"; + } + + OS << OptionalVersionEndif; + OS << OptionalExtensionEndif; + } + + OS << "\n// Disable any extensions we may have enabled previously.\n" + "#pragma OPENCL EXTENSION all : disable"; +} + void clang::EmitClangOpenCLBuiltins(RecordKeeper &Records, raw_ostream &OS) { BuiltinNameEmitter NameChecker(Records, OS); NameChecker.Emit(); } +void clang::EmitClangOpenCLBuiltinHeader(RecordKeeper &Records, + raw_ostream &OS) { + OpenCLBuiltinHeaderEmitter HeaderFileGenerator(Records, OS); + HeaderFileGenerator.emit(); +} + void clang::EmitClangOpenCLBuiltinTests(RecordKeeper &Records, raw_ostream &OS) { OpenCLBuiltinTestEmitter TestFileGenerator(Records, OS); Index: clang/utils/TableGen/TableGen.cpp =================================================================== --- clang/utils/TableGen/TableGen.cpp +++ clang/utils/TableGen/TableGen.cpp @@ -63,6 +63,7 @@ GenClangCommentCommandInfo, GenClangCommentCommandList, GenClangOpenCLBuiltins, + GenClangOpenCLBuiltinHeader, GenClangOpenCLBuiltinTests, GenArmNeon, GenArmFP16, @@ -195,6 +196,9 @@ "documentation comments"), clEnumValN(GenClangOpenCLBuiltins, "gen-clang-opencl-builtins", "Generate OpenCL builtin declaration handlers"), + clEnumValN(GenClangOpenCLBuiltinHeader, + "gen-clang-opencl-builtin-header", + "Generate OpenCL builtin header"), clEnumValN(GenClangOpenCLBuiltinTests, "gen-clang-opencl-builtin-tests", "Generate OpenCL builtin declaration tests"), clEnumValN(GenArmNeon, "gen-arm-neon", "Generate arm_neon.h for clang"), @@ -374,6 +378,9 @@ case GenClangOpenCLBuiltins: EmitClangOpenCLBuiltins(Records, OS); break; + case GenClangOpenCLBuiltinHeader: + EmitClangOpenCLBuiltinHeader(Records, OS); + break; case GenClangOpenCLBuiltinTests: EmitClangOpenCLBuiltinTests(Records, OS); break; Index: clang/utils/TableGen/TableGenBackends.h =================================================================== --- clang/utils/TableGen/TableGenBackends.h +++ clang/utils/TableGen/TableGenBackends.h @@ -122,6 +122,8 @@ void EmitClangOpenCLBuiltins(llvm::RecordKeeper &Records, llvm::raw_ostream &OS); +void EmitClangOpenCLBuiltinHeader(llvm::RecordKeeper &Records, + llvm::raw_ostream &OS); void EmitClangOpenCLBuiltinTests(llvm::RecordKeeper &Records, llvm::raw_ostream &OS);