[SYCL-Upstreaming] Add support for host kernel launch stmt generation (#51)

* Add support for host kernel launch stmt generation

This adds generation of a call to sycl_enqueue_kernel_launch function
aka "launcher" function. The launcher function can be a memeber of a
class or a free function defined at namespace scope. The lookup is
performed from SKEP attributed function scope. Because unqualified
lookup requires Scope object present and it only exists during parsing
stage and already EOLed at the point where templates instantiated, I had
to move some parts of SYCLKernelCallStmt generation to earlier stages
and now TreeTransform knows how to process SYCLKernelCallStmt.
I also had to invent a new expression - UnresolvedSYCLKernelExpr which
represents a string containing kernel name of a kernel that doesn't
exist yet. This expression is supposed to be transformed to a
StringLiteral during template instantiation phase. It should never reach
AST consumers like CodeGen of constexpr evaluators. This still requires
more testing and FIXME cleanups, but since it evolved into a quite
complicated patch I'm pushing it for earlier feedback.

* Remove a fixme from SemaSYCL

* Do not crash if original body was invalid

* Add AST test for skep-attributed member

* Fix a warning

* Extend codegen test a bit

* Find and replace

UnresolvedSYCLKernelNameExpr -> UnresolvedSYCLKernelLaunchExpr

* Implement the thing

* One more find and replace

* I don't know how it looks like

* Find and replace again

* Switch to UnresolvedSYCLKernelEntryPointStmt

* Apply suggestions from code review

* Remove log.txt

* Implement visiting

* Add tests

* Apply suggestions from code review

Co-authored-by: Tom Honermann <[email protected]>

* IdExpr -> KernelLaunchIdExpr

* Don't rely on compound

* UnresolvedSYCLKernelEntryPointStmt -> UnresolvedSYCLKernelCall

* Fix warnings

* Rename sycl_enqueue_kernel_launch -> sycl_kernel_launch

* Apply suggestions from code review

Co-authored-by: Tom Honermann <[email protected]>

* Remove array decay

* Add windows run line to the sema test

---------

Co-authored-by: Tom Honermann <[email protected]>

Author: Fznamznon

clang/include/clang/AST/RecursiveASTVisitor.h

@@ -2990,6 +2990,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

@@ -566,7 +566,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.
     }
@@ -634,7 +634,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.
 
@@ -702,7 +702,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:
 
@@ -713,16 +713,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
@@ -733,7 +733,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

@@ -13088,6 +13088,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

@@ -23,6 +23,7 @@ def CaseStmt : StmtNode<SwitchCase>;
 def DefaultStmt : StmtNode<SwitchCase>;
 def CapturedStmt : StmtNode<Stmt>;
 def SYCLKernelCallStmt : StmtNode<Stmt>;
+def UnresolvedSYCLKernelCallStmt : StmtNode<Stmt>;
 
 // Break/continue.
 def LoopControlStmt : StmtNode<Stmt, 1>;

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

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