[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

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";
 
for (const auto &Signature : FTypes) {
       OS << OptionalTypeExtEndif;
     }
 
     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; sort to follow order of definition in .td file.
+  std::vector<Record *> Builtins = Records.getAllDerivedDefinitions("Builtin");
+  llvm::sort(Builtins, LessRecord());
+
+  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) {
+      StringRef OptionalTypeExtEndif = emitTypeExtensionGuards(Signature);
+
+      // Emit function declaration.
+      OS << Signature[0] << " __ovld ";
+      if (B->getValueAsBit("IsConst"))
+        OS << "__cnfn ";
+      if (B->getValueAsBit("IsPure"))
+        OS << "__purefn ";
+      if (B->getValueAsBit("IsConv"))
+        OS << "__conv ";
+
+      OS << Name << "(";
+      if (Signature.size() > 1) {
+        for (unsigned I = 1; I < Signature.size(); I++) {
+          if (I != 1)
+            OS << ", ";
+          OS << Signature[I];
+        }
+      }
+      OS << ");\n";
+
+      OS << OptionalTypeExtEndif;
+    }
+
+    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