[SYCL] Add infrastructure for integration footer

clang/include/clang/Basic/LangOptions.h

@@ -350,6 +350,9 @@ class LangOptions : public LangOptionsBase {
   /// SYCL integration header to be generated by the device compiler
   std::string SYCLIntHeader;
 
+  /// SYCL integration footer to be generated by the device compiler
+  std::string SYCLIntFooter;
+
   LangOptions();
 
   // Define accessors/mutators for language options of enumeration type.

clang/include/clang/Driver/Options.td

@@ -5685,6 +5685,11 @@ def fsycl_int_header : Separate<["-"], "fsycl-int-header">,
   MarshallingInfoString<LangOpts<"SYCLIntHeader">>;
 def fsycl_int_header_EQ : Joined<["-"], "fsycl-int-header=">,
   Alias<fsycl_int_header>;
+def fsycl_int_footer : Separate<["-"], "fsycl-int-footer">,
+  HelpText<"Generate SYCL integration footer into this file.">,
+  MarshallingInfoString<LangOpts<"SYCLIntFooter">>;
+def fsycl_int_footer_EQ : Joined<["-"], "fsycl-int-footer=">,
+  Alias<fsycl_int_footer>;
 def fsycl_std_layout_kernel_params: Flag<["-"], "fsycl-std-layout-kernel-params">,
   HelpText<"Enable standard layout requirement for SYCL kernel parameters.">,
   MarshallingInfoFlag<LangOpts<"SYCLStdLayoutKernelParams">>;

clang/include/clang/Sema/Sema.h

@@ -320,15 +320,14 @@ class SYCLIntegrationHeader {
   };
 
 public:
-  SYCLIntegrationHeader(DiagnosticsEngine &Diag, bool UnnamedLambdaSupport,
-                        Sema &S);
+  SYCLIntegrationHeader(bool UnnamedLambdaSupport, Sema &S);
 
   /// Emits contents of the header into given stream.
   void emit(raw_ostream &Out);
 
   /// Emits contents of the header into a file with given name.
   /// Returns true/false on success/failure.
-  bool emit(const StringRef &MainSrc);
+  bool emit(StringRef MainSrc);
 
   ///  Signals that subsequent parameter descriptor additions will go to
   ///  the kernel with given name. Starts new kernel invocation descriptor.
@@ -431,6 +430,16 @@ class SYCLIntegrationHeader {
   Sema &S;
 };
 
+class SYCLIntegrationFooter {
+public:
+  SYCLIntegrationFooter(Sema &S) : S(S) {}
+  bool emit(StringRef MainSrc);
+
+private:
+  bool emit(raw_ostream &O);
+  Sema &S;
+};
+
 /// Tracks expected type during expression parsing, for use in code completion.
 /// The type is tied to a particular token, all functions that update or consume
 /// the type take a start location of the token they are looking at as a
@@ -13110,6 +13119,7 @@ class Sema final {
   // SYCL integration header instance for current compilation unit this Sema
   // is associated with.
   std::unique_ptr<SYCLIntegrationHeader> SyclIntHeader;
+  std::unique_ptr<SYCLIntegrationFooter> SyclIntFooter;
 
   // Used to suppress diagnostics during kernel construction, since these were
   // already emitted earlier. Diagnosing during Kernel emissions also skips the
@@ -13124,10 +13134,16 @@ class Sema final {
   SYCLIntegrationHeader &getSyclIntegrationHeader() {
     if (SyclIntHeader == nullptr)
       SyclIntHeader = std::make_unique<SYCLIntegrationHeader>(
-          getDiagnostics(), getLangOpts().SYCLUnnamedLambda, *this);
+          getLangOpts().SYCLUnnamedLambda, *this);
     return *SyclIntHeader.get();
   }
 
+  SYCLIntegrationFooter &getSyclIntegrationFooter() {
+    if (SyclIntFooter == nullptr)
+      SyclIntFooter = std::make_unique<SYCLIntegrationFooter>(*this);
+    return *SyclIntFooter.get();
+  }
+
   enum SYCLRestrictKind {
     KernelGlobalVariable,
     KernelRTTI,

clang/lib/Frontend/CompilerInvocation.cpp

@@ -3698,8 +3698,8 @@ bool CompilerInvocation::ParseLangArgs(LangOptions &Opts, ArgList &Args,
   // '-mignore-xcoff-visibility' is implied. The generated command line will
   // contain both '-fvisibility default' and '-mignore-xcoff-visibility' and
   // subsequent calls to `CreateFromArgs`/`generateCC1CommandLine` will always
-  // produce the same arguments. 
- 
+  // produce the same arguments.
+
   if (T.isOSAIX() && (Args.hasArg(OPT_mignore_xcoff_visibility) ||
                       !Args.hasArg(OPT_fvisibility)))
     Opts.IgnoreXCOFFVisibility = 1;
@@ -4291,6 +4291,22 @@ static bool ParseTargetArgs(TargetOptions &Opts, ArgList &Args,
   return Success && Diags.getNumErrors() == NumErrorsBefore;
 }
 
+static void CreateEmptyFile(StringRef HeaderName) {
+  if (HeaderName.empty())
+    return;
+
+  Expected<llvm::sys::fs::file_t> FT = llvm::sys::fs::openNativeFileForWrite(
+      HeaderName, llvm::sys::fs::CD_OpenAlways, llvm::sys::fs::OF_None);
+  if (FT)
+    llvm::sys::fs::closeFile(*FT);
+  else {
+    // Emit a message but don't terminate; compilation will fail
+    // later if this file is absent.
+    llvm::errs() << "Error: " << llvm::toString(FT.takeError())
+                 << " when opening " << HeaderName << "\n";
+  }
+}
+
 bool CompilerInvocation::CreateFromArgsImpl(
     CompilerInvocation &Res, ArrayRef<const char *> CommandLineArgs,
     DiagnosticsEngine &Diags, const char *Argv0) {
@@ -4372,21 +4388,9 @@ bool CompilerInvocation::CreateFromArgsImpl(
   if (LangOpts.SYCLIsDevice) {
     // Set the triple of the host for SYCL device compile.
     Res.getTargetOpts().HostTriple = Res.getFrontendOpts().AuxTriple;
-    // If specified, create an empty integration header file for now.
-    const StringRef &HeaderName = LangOpts.SYCLIntHeader;
-    if (!HeaderName.empty()) {
-      Expected<llvm::sys::fs::file_t> ft =
-          llvm::sys::fs::openNativeFileForWrite(
-              HeaderName, llvm::sys::fs::CD_OpenAlways, llvm::sys::fs::OF_None);
-      if (ft)
-        llvm::sys::fs::closeFile(*ft);
-      else {
-        // Emit a message but don't terminate; compilation will fail
-        // later if this file is absent.
-        llvm::errs() << "Error: " << llvm::toString(ft.takeError())
-                     << " when opening " << HeaderName << "\n";
-      }
-    }
+    // If specified, create empty integration header files for now.
+    CreateEmptyFile(LangOpts.SYCLIntHeader);
+    CreateEmptyFile(LangOpts.SYCLIntFooter);
   }
 
   Success &= ParseCodeGenArgs(Res.getCodeGenOpts(), Args, DashX, Diags, T,

clang/lib/Sema/Sema.cpp

@@ -183,7 +183,7 @@ Sema::Sema(Preprocessor &pp, ASTContext &ctxt, ASTConsumer &consumer,
       DisableTypoCorrection(false), TyposCorrected(0), AnalysisWarnings(*this),
       ThreadSafetyDeclCache(nullptr), VarDataSharingAttributesStack(nullptr),
       CurScope(nullptr), Ident_super(nullptr), Ident___float128(nullptr),
-      SyclIntHeader(nullptr) {
+      SyclIntHeader(nullptr), SyclIntFooter(nullptr) {
   TUScope = nullptr;
   isConstantEvaluatedOverride = false;
 
@@ -1037,6 +1037,8 @@ void Sema::ActOnEndOfTranslationUnitFragment(TUFragmentKind Kind) {
     // Emit SYCL integration header for current translation unit if needed
     if (SyclIntHeader != nullptr)
       SyclIntHeader->emit(getLangOpts().SYCLIntHeader);
+    if (SyclIntFooter != nullptr)
+      SyclIntFooter->emit(getLangOpts().SYCLIntFooter);
     MarkDevice();
   }
 

clang/lib/Sema/SemaSYCL.cpp

@@ -3086,6 +3086,14 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
   using SyclKernelFieldHandler::leaveStruct;
 };
 
+class SyclKernelIntFooterCreator : public SyclKernelFieldHandler {
+  SYCLIntegrationFooter &Footer;
+
+public:
+  SyclKernelIntFooterCreator(Sema &S, SYCLIntegrationFooter &F)
+      : SyclKernelFieldHandler(S), Footer(F) {}
+};
+
 } // namespace
 
 class SYCLKernelNameTypeVisitor
@@ -3411,9 +3419,13 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc,
       calculateKernelNameType(Context, KernelCallerFunc), KernelName,
       StableName, KernelCallerFunc);
 
+  SyclKernelIntFooterCreator int_footer(*this, getSyclIntegrationFooter());
+
   KernelObjVisitor Visitor{*this};
-  Visitor.VisitRecordBases(KernelObj, kernel_decl, kernel_body, int_header);
-  Visitor.VisitRecordFields(KernelObj, kernel_decl, kernel_body, int_header);
+  Visitor.VisitRecordBases(KernelObj, kernel_decl, kernel_body, int_header,
+                           int_footer);
+  Visitor.VisitRecordFields(KernelObj, kernel_decl, kernel_body, int_header,
+                            int_footer);
 
   if (ParmVarDecl *KernelHandlerArg =
           getSyclKernelHandlerArg(KernelCallerFunc)) {
@@ -4149,7 +4161,7 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
   O << "\n";
 }
 
-bool SYCLIntegrationHeader::emit(const StringRef &IntHeaderName) {
+bool SYCLIntegrationHeader::emit(StringRef IntHeaderName) {
   if (IntHeaderName.empty())
     return false;
   int IntHeaderFD = 0;
@@ -4221,11 +4233,31 @@ void SYCLIntegrationHeader::setCallsThisGroup(bool B) {
   K->FreeFunctionCalls.CallsThisGroup = B;
 }
 
-SYCLIntegrationHeader::SYCLIntegrationHeader(DiagnosticsEngine &_Diag,
-                                             bool _UnnamedLambdaSupport,
+SYCLIntegrationHeader::SYCLIntegrationHeader(bool _UnnamedLambdaSupport,
                                              Sema &_S)
     : UnnamedLambdaSupport(_UnnamedLambdaSupport), S(_S) {}
 
+// Post-compile integration header support.
+bool SYCLIntegrationFooter::emit(StringRef IntHeaderName) {
+  if (IntHeaderName.empty())
+    return false;
+  int IntHeaderFD = 0;
+  std::error_code EC =
+      llvm::sys::fs::openFileForWrite(IntHeaderName, IntHeaderFD);
+  if (EC) {
+    llvm::errs() << "Error: " << EC.message() << "\n";
+    // compilation will fail on absent include file - don't need to fail here
+    return false;
+  }
+  llvm::raw_fd_ostream Out(IntHeaderFD, true /*close in destructor*/);
+  return emit(Out);
+}
+
+bool SYCLIntegrationFooter::emit(raw_ostream &O) {
+  O << "// Integration Footer contents to go here.\n";
+  return true;
+}
+
 // -----------------------------------------------------------------------------
 // Utility class methods
 // -----------------------------------------------------------------------------

clang/test/CodeGenSYCL/integration_footer.cpp

@@ -0,0 +1,10 @@
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-footer=%t.h %s -emit-llvm -o %t.ll
+// RUN: FileCheck -input-file=%t.h %s
+
+// CHECK: // Integration Footer contents to go here.
+
+#include "Inputs/sycl.hpp"
+
+int main() {
+  cl::sycl::kernel_single_task<class first_kernel>([]() {});
+}