[OpenCL] Add TableGen emitter for OpenCL builtin header

[svenvh]

Summary

Add an emitter to produce something similar to opencl-c.h from the
OpenCL builtin descriptions in OpenCLBuiltins.td

Factor out functionality that can be shared between the test emitter
and the header emitter into a new base class.

This only adds the emitter, without any use of it currently.

Example output excerpt (the full output is about 16k lines):

...
__ovld __cnfn size_t get_global_id(uint);
__ovld __cnfn size_t get_local_size(uint);
__ovld __cnfn size_t get_local_id(uint);
__ovld __cnfn size_t get_num_groups(uint);
__ovld __cnfn size_t get_group_id(uint);
__ovld __cnfn size_t get_global_offset(uint);
#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
__ovld size_t get_enqueued_local_size(uint);
#endif // MinVersion
...

[azabaznov]

Thanks for working on this! That's cool, I managed to reproduce something similar to https://reviews.llvm.org/D99577.

[Anastasia]

Do you think we could add any test for this yet? Or do we have to wait until we have other pieces in place for diffing to opencl-c.h? We could add a test for parsing the header but it is probably going to be very slow test?

[svenvh]

In D104040#2824281, @Anastasia wrote:

Do you think we could add any test for this yet? Or do we have to wait until we have other pieces in place for diffing to opencl-c.h? We could add a test for parsing the header but it is probably going to be very slow test?

The question is what we actually want to test at this point? I am wondering if it would make more sense to work towards equivalence of opencl-c.h and the generated header file. We can prove equivalence offline, as a one-off without needing to commit changes to the AST printer. If they are equivalent, we can commit the generated header file in place of opencl-c.h. After that, adding a test for diffing both should be fairly trivial.

[svenvh]

Rebased on latest main.

Also takes TypeExtensions into account now.

clang/utils/TableGen/ClangOpenCLBuiltinEmitter.cpp

private:
   // SignatureListMap[<I1, I2, I3>] = <
   //                  <"cos", "sin">,
   //                  <Signature1, Signature2, Signature3>>
   // The function "tan", having the same signatures, would be mapped to the
   // same entry (<I1, I2, I3>).
   MapVector<BuiltinIndexListTy *, BuiltinTableEntries> 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;
     bool IsVolatile : 1;
     bool IsPointer : 1;
     StringRef AddrSpace;
   };
 
protected:
   //   [char, float, float]
   //   [char, float2, float2]
   //   [char, float3, float3]
   //   ...
   void
   expandTypesInSignature(const std::vector<Record *> &Signature,
                          SmallVectorImpl<SmallVector<std::string, 2>> &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;
 
   // The output file.
   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() {
   emitSourceFileHeader("OpenCL Builtin handling", OS);
 
   OS << "#include \"llvm/ADT/StringRef.h\"\n";
   OS << "using namespace clang;\n\n";
 
if (Ty.IsPointer != 0) {
     }
   }
 )";
 
   // End of the "OCL2Qual" function.
   OS << "\n} // OCL2Qual\n";
 }
 
-std::string OpenCLBuiltinTestEmitter::getTypeString(const Record *Type,
+std::string OpenCLBuiltinFileEmitterBase::getTypeString(const Record *Type,
                                                         TypeFlags Flags,
                                                         int VectorSize) const {
   std::string S;
   if (Type->getValueAsBit("IsConst") || Flags.IsConst) {
     S += "const ";
   }
   if (Type->getValueAsBit("IsVolatile") || Flags.IsVolatile) {
     S += "volatile ";
   }
 
std::string OpenCLBuiltinFileEmitterBase::getTypeString(const Record *Type,
 
   if (Type->getValueAsBit("IsPointer") || Flags.IsPointer) {
     S += " *";
   }
 
   return S;
 }
 
-void OpenCLBuiltinTestEmitter::getTypeLists(
+void OpenCLBuiltinFileEmitterBase::getTypeLists(
     Record *Type, TypeFlags &Flags, std::vector<Record *> &TypeList,
     std::vector<int64_t> &VectorList) const {
   bool isGenType = Type->isSubClassOf("GenericType");
   if (isGenType) {
     TypeList = Type->getValueAsDef("TypeList")->getValueAsListOfDefs("List");
     VectorList =
         Type->getValueAsDef("VectorList")->getValueAsListOfInts("List");
     return;
if (Type->isSubClassOf("PointerType") || Type->isSubClassOf("ConstType") ||
     }
   }
 
   // Not a GenericType, so just insert the single type.
   TypeList.push_back(Type);
   VectorList.push_back(Type->getValueAsInt("VecWidth"));
 }
 
-void OpenCLBuiltinTestEmitter::expandTypesInSignature(
+void OpenCLBuiltinFileEmitterBase::expandTypesInSignature(
     const std::vector<Record *> &Signature,
     SmallVectorImpl<SmallVector<std::string, 2>> &Types) {
   // Find out if there are any GenTypes in this signature, and if so, calculate
   // into how many signatures they will expand.
   unsigned NumSignatures = 1;
   SmallVector<SmallVector<std::string, 4>, 4> ExpandedGenTypes;
   for (const auto &Arg : Signature) {
     SmallVector<std::string, 4> ExpandedArg;
for (unsigned ArgNum = 0; ArgNum < Signature.size(); ArgNum++) {
       // GenTypes just repeat, so index modulo the number of expanded types.
       size_t TypeIndex = I % ExpandedGenTypes[ArgNum].size();
       Args.push_back(ExpandedGenTypes[ArgNum][TypeIndex]);
     }
     Types.push_back(Args);
   }
 }
 
-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
 #pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
 #pragma OPENCL EXTENSION cl_khr_int64_extended_atomics : enable
 #pragma OPENCL EXTENSION cl_khr_gl_msaa_sharing : enable
 #pragma OPENCL EXTENSION cl_khr_mipmap_image_writes : enable
 #pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable
 
 )";
+}
 
-  // Ensure each test has a unique name by numbering them.
-  unsigned TestID = 0;
-
-  // Iterate over all builtins.
-  std::vector<Record *> Builtins = Records.getAllDerivedDefinitions("Builtin");
-  for (const auto *B : Builtins) {
-    StringRef Name = B->getValueAsString("Name");
-
-    SmallVector<SmallVector<std::string, 2>, 4> FTypes;
-    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";
+std::string
+OpenCLBuiltinFileEmitterBase::emitExtensionGuard(const Record *Builtin) {
+      Builtin->getValueAsDef("Extension")->getValueAsString("ExtName");
+  if (Extensions.empty())
+    return "";
+
 
-      SmallVector<StringRef, 2> 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;
-    }
+
+  return "#endif // Extension\n";
+
+std::string
+OpenCLBuiltinFileEmitterBase::emitVersionGuard(const Record *Builtin) {
+  std::string OptionalEndif;
+  int MinVersion = Builtin->getValueAsDef("MinVersion")->getValueAsInt("ID");
+  int MaxVersion = Builtin->getValueAsDef("MaxVersion")->getValueAsInt("ID");
+  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;
+
+  // Iterate over all builtins.
+  std::vector<Record *> Builtins = Records.getAllDerivedDefinitions("Builtin");
+  for (const auto *B : Builtins) {
+    StringRef Name = B->getValueAsString("Name");
+
+    SmallVector<SmallVector<std::string, 2>, 4> FTypes;
+    expandTypesInSignature(B->getValueAsListOfDefs("Signature"), FTypes);
+
+    OS << "// Test " << Name << "\n";
+
+    std::string OptionalExtensionEndif = emitExtensionGuard(B);
+    std::string OptionalVersionEndif = emitVersionGuard(B);
+
     for (const auto &Signature : FTypes) {
       // Emit function declaration.
       OS << Signature[0] << " test" << TestID++ << "_" << Name << "(";
       if (Signature.size() > 1) {
         for (unsigned I = 1; I < Signature.size(); I++) {
           if (I != 1)
             OS << ", ";
           OS << Signature[I] << " arg" << I;
for (const auto &Signature : FTypes) {
           OS << ", ";
         OS << "arg" << I;
       }
       OS << ");\n";
 
       // 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<Record *> 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<SmallVector<std::string, 2>, 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);
   TestFileGenerator.emit();
 }

clang/utils/TableGen/TableGen.cpp

enum ActionType {
   GenClangSyntaxNodeList,
   GenClangSyntaxNodeClasses,
   GenClangCommentHTMLTags,
   GenClangCommentHTMLTagsProperties,
   GenClangCommentHTMLNamedCharacterReferences,
   GenClangCommentCommandInfo,
   GenClangCommentCommandList,
   GenClangOpenCLBuiltins,
+  GenClangOpenCLBuiltinHeader,
   GenClangOpenCLBuiltinTests,
   GenArmNeon,
   GenArmFP16,
   GenArmBF16,
   GenArmNeonSema,
   GenArmNeonTest,
   GenArmMveHeader,
   GenArmMveBuiltinDef,
cl::values(
         clEnumValN(GenClangCommentCommandInfo, "gen-clang-comment-command-info",
                    "Generate command properties for commands that "
                    "are used in documentation comments"),
         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(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"),
         clEnumValN(GenArmFP16, "gen-arm-fp16", "Generate arm_fp16.h for clang"),
         clEnumValN(GenArmBF16, "gen-arm-bf16", "Generate arm_bf16.h for clang"),
         clEnumValN(GenArmNeonSema, "gen-arm-neon-sema",
                    "Generate ARM NEON sema support for clang"),
         clEnumValN(GenArmNeonTest, "gen-arm-neon-test",
case GenClangCommentCommandInfo:
     EmitClangCommentCommandInfo(Records, OS);
     break;
   case GenClangCommentCommandList:
     EmitClangCommentCommandList(Records, OS);
     break;
   case GenClangOpenCLBuiltins:
     EmitClangOpenCLBuiltins(Records, OS);
     break;
+  case GenClangOpenCLBuiltinHeader:
+    EmitClangOpenCLBuiltinHeader(Records, OS);
+    break;
   case GenClangOpenCLBuiltinTests:
     EmitClangOpenCLBuiltinTests(Records, OS);
     break;
   case GenClangSyntaxNodeList:
     EmitClangSyntaxNodeList(Records, OS);
     break;
   case GenClangSyntaxNodeClasses:
     EmitClangSyntaxNodeClasses(Records, OS);

clang/utils/TableGen/TableGenBackends.h

 void EmitCdeBuiltinAliases(llvm::RecordKeeper &Records, llvm::raw_ostream &OS);
 
 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 EmitClangOpenCLBuiltinHeader(llvm::RecordKeeper &Records,
+                                  llvm::raw_ostream &OS);
 void EmitClangOpenCLBuiltinTests(llvm::RecordKeeper &Records,
                                  llvm::raw_ostream &OS);
 
 void EmitClangDataCollectors(llvm::RecordKeeper &Records,
                              llvm::raw_ostream &OS);
 
 void EmitTestPragmaAttributeSupportedAttributes(llvm::RecordKeeper &Records,
                                                 llvm::raw_ostream &OS);
 
 } // end namespace clang
 
 #endif