[SYCL-Upstreaming] Add support for host kernel launch stmt generation

clang/include/clang/AST/RecursiveASTVisitor.h

@@ -2999,6 +2999,13 @@ DEF_TRAVERSE_STMT(ParenListExpr, {})
 DEF_TRAVERSE_STMT(SYCLUniqueStableNameExpr, {
   TRY_TO(TraverseTypeLoc(S->getTypeSourceInfo()->getTypeLoc()));
 })
+DEF_TRAVERSE_STMT(UnresolvedSYCLKernelCallStmt, {
+  if (getDerived().shouldVisitImplicitCode()) {
+    TRY_TO(TraverseStmt(S->getOriginalStmt()));
+    TRY_TO(TraverseStmt(S->getKernelLaunchIdExpr()));
+    ShouldVisitChildren = false;
+  }
+})
 DEF_TRAVERSE_STMT(OpenACCAsteriskSizeExpr, {})
 DEF_TRAVERSE_STMT(PredefinedExpr, {})
 DEF_TRAVERSE_STMT(ShuffleVectorExpr, {})

clang/include/clang/AST/StmtSYCL.h

@@ -99,6 +99,59 @@ class SYCLKernelCallStmt : public Stmt {
   }
 };
 
+// UnresolvedSYCLKernelCallStmt represents an invocation of a SYCL kernel in
+// a dependent context for which lookup of the sycl_enqueue_kernel_launch
+// identifier cannot be performed. These statements are transformed to
+// SYCLKernelCallStmt during template instantiation.
+class UnresolvedSYCLKernelCallStmt : public Stmt {
+  friend class ASTStmtReader;
+  Stmt *OriginalStmt = nullptr;
+  // KernelLaunchIdExpr stores an UnresolvedLookupExpr or UnresolvedMemberExpr
+  // corresponding to the SYCL kernel launch function for which a call
+  // will be synthesized during template instantiation.
+  Expr *KernelLaunchIdExpr = nullptr;
+  UnresolvedSYCLKernelCallStmt(CompoundStmt *CS, Expr *IdExpr)
+      : Stmt(UnresolvedSYCLKernelCallStmtClass), OriginalStmt(CS),
+        KernelLaunchIdExpr(IdExpr) {}
+
+  void setKernelLaunchIdExpr(Expr *IdExpr) { KernelLaunchIdExpr = IdExpr; }
+  void setOriginalStmt(CompoundStmt *CS) { OriginalStmt = CS; }
+
+public:
+  static UnresolvedSYCLKernelCallStmt *
+  Create(const ASTContext &C, CompoundStmt *CS, Expr *IdExpr) {
+    return new (C) UnresolvedSYCLKernelCallStmt(CS, IdExpr);
+  }
+
+  static UnresolvedSYCLKernelCallStmt *CreateEmpty(const ASTContext &C) {
+    return new (C) UnresolvedSYCLKernelCallStmt(nullptr, nullptr);
+  }
+
+  Expr *getKernelLaunchIdExpr() const { return KernelLaunchIdExpr; }
+  CompoundStmt *getOriginalStmt() { return cast<CompoundStmt>(OriginalStmt); }
+  const CompoundStmt *getOriginalStmt() const {
+    return cast<CompoundStmt>(OriginalStmt);
+  }
+
+  SourceLocation getBeginLoc() const LLVM_READONLY {
+    return getOriginalStmt()->getBeginLoc();
+  }
+
+  SourceLocation getEndLoc() const LLVM_READONLY {
+    return getOriginalStmt()->getEndLoc();
+  }
+  static bool classof(const Stmt *T) {
+    return T->getStmtClass() == UnresolvedSYCLKernelCallStmtClass;
+  }
+  child_range children() {
+    return child_range(&OriginalStmt, &OriginalStmt + 1);
+  }
+
+  const_child_range children() const {
+    return const_child_range(&OriginalStmt, &OriginalStmt + 1);
+  }
+};
+
 } // end namespace clang
 
 #endif // LLVM_CLANG_AST_STMTSYCL_H

clang/include/clang/Basic/AttrDocs.td

@@ -554,7 +554,7 @@ follows.
   namespace sycl {
   class handler {
     template<typename KernelNameType, typename... Ts>
-    void sycl_enqueue_kernel_launch(const char *KernelName, Ts...) {
+    void sycl_kernel_launch(const char *KernelName, Ts...) {
       // Call functions appropriate for the desired offload backend
       // (OpenCL, CUDA, HIP, Level Zero, etc...) to enqueue kernel invocation.
     }
@@ -622,7 +622,7 @@ The offload kernel entry point for a SYCL kernel performs the following tasks:
 The ``sycl_kernel_entry_point`` attribute facilitates or automates these tasks
 by generating the offload kernel entry point, generating a unique symbol name
 for it, synthesizing code for kernel argument decomposition and reconstruction,
-and synthesizing a call to a ``sycl_enqueue_kernel_launch`` function template
+and synthesizing a call to a ``sycl_kernel_launch`` function template
 with the kernel name type, kernel symbol name, and (decomposed) kernel arguments
 passed as template or function arguments.
 
@@ -690,7 +690,7 @@ replaced with synthesized code that looks approximately as follows.
 
   sycl::stream sout = Kernel.sout;
   S s = Kernel.s;
-  sycl_enqueue_kernel_launch<KN>("kernel-symbol-name", sout, s);
+  sycl_kernel_launch<KN>("kernel-symbol-name", sout, s);
 
 There are a few items worthy of note:
 
@@ -701,16 +701,16 @@ There are a few items worthy of note:
 #. ``kernel-symbol-name`` is substituted for the actual symbol name that would
    be generated; these names are implementation details subject to change.
 
-#. Lookup for the ``sycl_enqueue_kernel_launch()`` function template is
+#. Lookup for the ``sycl_kernel_launch()`` function template is
    performed from the (possibly instantiated) location of the definition of
    ``kernel_entry_point()``. If overload resolution fails, the program is
    ill-formed. If the selected overload is a non-static member function, then
    ``this`` is passed for the implicit object parameter.
 
-#. Function arguments passed to ``sycl_enqueue_kernel_launch()`` are passed
+#. Function arguments passed to ``sycl_kernel_launch()`` are passed
    as if by ``std::forward<X>(x)``.
 
-#. The ``sycl_enqueue_kernel_launch()`` function is expected to be provided by
+#. The ``sycl_kernel_launch()`` function is expected to be provided by
    the SYCL library implementation. It is responsible for scheduling execution
    of the generated offload kernel entry point identified by
    ``kernel-symbol-name`` and copying the (decomposed) kernel arguments to
@@ -721,7 +721,7 @@ attribute to be called for the offload kernel entry point to be emitted. For
 inline functions and function templates, any ODR-use will suffice. For other
 functions, an ODR-use is not required; the offload kernel entry point will be
 emitted if the function is defined. In any case, a call to the function is
-required for the synthesized call to ``sycl_enqueue_kernel_launch()`` to occur.
+required for the synthesized call to ``sycl_kernel_launch()`` to occur.
 
 Functions declared with the ``sycl_kernel_entry_point`` attribute are not
 limited to the simple example shown above. They may have additional template

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -13015,6 +13015,15 @@ def err_sycl_entry_point_return_type : Error<
 def err_sycl_entry_point_deduced_return_type : Error<
   "the %0 attribute only applies to functions with a non-deduced 'void' return"
   " type">;
+def err_sycl_host_no_launch_function : Error<
+  "unable to find suitable 'sycl_kernel_launch' function for host code "
+  "synthesis">;
+def warn_sycl_device_no_host_launch_function : Warning<
+  "unable to find suitable 'sycl_kernel_launch' function for host code "
+  "synthesis">,
+  InGroup<DiagGroup<"sycl-host-launcher">>;
+def note_sycl_host_launch_function : Note<
+  "define 'sycl_kernel_launch' function template to fix this problem">;
 
 def warn_cuda_maxclusterrank_sm_90 : Warning<
   "maxclusterrank requires sm_90 or higher, CUDA arch provided: %0, ignoring "

clang/include/clang/Basic/StmtNodes.td

@@ -25,6 +25,7 @@ def CaseStmt : StmtNode<SwitchCase>;
 def DefaultStmt : StmtNode<SwitchCase>;
 def CapturedStmt : StmtNode<Stmt>;
 def SYCLKernelCallStmt : StmtNode<Stmt>;
+def UnresolvedSYCLKernelCallStmt : StmtNode<Stmt>;
 
 // Statements that might produce a value (for example, as the last non-null
 // statement in a GNU statement-expression).

clang/include/clang/Sema/ScopeInfo.h

@@ -245,6 +245,10 @@ class FunctionScopeInfo {
   /// The set of GNU address of label extension "&&label".
   llvm::SmallVector<AddrLabelExpr *, 4> AddrLabels;
 
+  /// An unresolved identifier lookup expression for an implicit call
+  /// to a SYCL kernel launch function in a dependent context.
+  Expr *SYCLKernelLaunchIdExpr = nullptr;
+
 public:
   /// Represents a simple identification of a weak object.
   ///

clang/include/clang/Sema/SemaSYCL.h

@@ -66,7 +66,11 @@ class SemaSYCL : public SemaBase {
 
   void CheckSYCLExternalFunctionDecl(FunctionDecl *FD);
   void CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD);
-  StmtResult BuildSYCLKernelCallStmt(FunctionDecl *FD, CompoundStmt *Body);
+  StmtResult BuildSYCLKernelCallStmt(FunctionDecl *FD, CompoundStmt *Body,
+                                     Expr *LaunchIdExpr);
+  ExprResult BuildSYCLKernelLaunchIdExpr(FunctionDecl *FD, QualType KNT);
+  StmtResult BuildUnresolvedSYCLKernelCallStmt(CompoundStmt *CS,
+                                                     Expr *IdExpr);
 };
 
 } // namespace clang

clang/include/clang/Serialization/ASTBitCodes.h

@@ -1615,6 +1615,9 @@ enum StmtCode {
   /// A SYCLKernelCallStmt record.
   STMT_SYCLKERNELCALL,
 
+  /// A SYCLKernelCallStmt record.
+  STMT_UNRESOLVED_SYCL_KERNEL_CALL,
+
   /// A GCC-style AsmStmt record.
   STMT_GCCASM,
 

clang/lib/AST/ComputeDependence.cpp

@@ -16,6 +16,7 @@
 #include "clang/AST/ExprConcepts.h"
 #include "clang/AST/ExprObjC.h"
 #include "clang/AST/ExprOpenMP.h"
+#include "clang/AST/StmtSYCL.h"
 #include "clang/Basic/ExceptionSpecificationType.h"
 #include "llvm/ADT/ArrayRef.h"
 

clang/lib/AST/StmtPrinter.cpp

@@ -1428,6 +1428,11 @@ void StmtPrinter::VisitSYCLUniqueStableNameExpr(
   OS << ")";
 }
 
+void StmtPrinter::VisitUnresolvedSYCLKernelCallStmt(
+    UnresolvedSYCLKernelCallStmt *Node) {
+  PrintStmt(Node->getOriginalStmt());
+}
+
 void StmtPrinter::VisitPredefinedExpr(PredefinedExpr *Node) {
   OS << PredefinedExpr::getIdentKindName(Node->getIdentKind());
 }

clang/lib/AST/StmtProfile.cpp

@@ -1374,6 +1374,11 @@ void StmtProfiler::VisitSYCLUniqueStableNameExpr(
   VisitType(S->getTypeSourceInfo()->getType());
 }
 
+void StmtProfiler::VisitUnresolvedSYCLKernelCallStmt(
+    const UnresolvedSYCLKernelCallStmt *S) {
+  VisitStmt(S);
+}
+
 void StmtProfiler::VisitPredefinedExpr(const PredefinedExpr *S) {
   VisitExpr(S);
   ID.AddInteger(llvm::to_underlying(S->getIdentKind()));

clang/lib/CodeGen/CGStmt.cpp

@@ -117,6 +117,7 @@ void CodeGenFunction::EmitStmt(const Stmt *S, ArrayRef<const Attr *> Attrs) {
   case Stmt::CaseStmtClass:
   case Stmt::SEHLeaveStmtClass:
   case Stmt::SYCLKernelCallStmtClass:
+  case Stmt::UnresolvedSYCLKernelCallStmtClass:
     llvm_unreachable("should have emitted these statements as simple");
 
 #define STMT(Type, Base)

clang/lib/Sema/SemaDecl.cpp

@@ -15754,7 +15754,6 @@ Sema::ActOnStartOfFunctionDef(Scope *FnBodyScope, Declarator &D,
   if (!Bases.empty())
     OpenMP().ActOnFinishedFunctionDefinitionInOpenMPDeclareVariantScope(Dcl,
                                                                         Bases);
-
   return Dcl;
 }
 
@@ -16167,6 +16166,20 @@ Decl *Sema::ActOnStartOfFunctionDef(Scope *FnBodyScope, Decl *D,
 
   maybeAddDeclWithEffects(FD);
 
+  if (FD && !FD->isInvalidDecl() &&
+      FD->hasAttr<SYCLKernelEntryPointAttr>() && FnBodyScope) {
+    // Building KernelLaunchIdExpr requires performing an unqualified lookup
+    // which can only be done correctly while the stack of parsing scopes is
+    // alive, so we do it here when we start parsing function body even if it is
+    // a templated function.
+    const auto *SKEPAttr = FD->getAttr<SYCLKernelEntryPointAttr>();
+    if (!SKEPAttr->isInvalidAttr()) {
+      ExprResult LaunchIdExpr =
+          SYCL().BuildSYCLKernelLaunchIdExpr(FD, SKEPAttr->getKernelName());
+      getCurFunction()->SYCLKernelLaunchIdExpr = LaunchIdExpr.get();
+    }
+  }
+
   return D;
 }
 
@@ -16368,9 +16381,21 @@ Decl *Sema::ActOnFinishFunctionBody(Decl *dcl, Stmt *Body, bool IsInstantiation,
       SKEPAttr->setInvalidAttr();
     }
 
-    if (Body && !FD->isTemplated() && !SKEPAttr->isInvalidAttr()) {
-      StmtResult SR =
-          SYCL().BuildSYCLKernelCallStmt(FD, cast<CompoundStmt>(Body));
+    // We don't need to build SYCLKernelCallStmt for template instantiations
+    // since it was already created by template instantiator.
+    if (Body && !SKEPAttr->isInvalidAttr()) {
+      StmtResult SR;
+      if (FD->isTemplated()) {
+        SR = SYCL().BuildUnresolvedSYCLKernelCallStmt(
+            cast<CompoundStmt>(Body), getCurFunction()->SYCLKernelLaunchIdExpr);
+      } else if (FD->isTemplateInstantiation()) {
+        assert(isa<SYCLKernelCallStmt>(Body));
+        SR = Body;
+      } else {
+        SR = SYCL().BuildSYCLKernelCallStmt(
+            FD, cast<CompoundStmt>(Body),
+            getCurFunction()->SYCLKernelLaunchIdExpr);
+      }
       if (SR.isInvalid())
         return nullptr;
       Body = SR.get();

clang/lib/Sema/SemaExceptionSpec.cpp

@@ -1261,6 +1261,9 @@ CanThrowResult Sema::canThrow(const Stmt *S) {
     return canSubStmtsThrow(*this, SKCS->getKernelLaunchStmt());
   }
 
+  case Stmt::UnresolvedSYCLKernelCallStmtClass:
+    return CT_Dependent;
+
     // ObjC message sends are like function calls, but never have exception
     // specs.
   case Expr::ObjCMessageExprClass: