[SYCL] SYCL host kernel launch support for the sycl_kernel_entry_point attribute.

The `sycl_kernel_entry_point` attribute facilitates the generation of an
offload kernel entry point function with parameters corresponding to the
(potentially decomposed) kernel arguments and a body that (potentially
reconstructs the arguments and) executes the kernel. This change adds
symmetric support for the SYCL host through an interface that provides
symbol names and (potentially decomposed) kernel arguments to the SYCL
library.

Consider the following function declared with the `sycl_kernel_entry_point`
attribute with a call to this function occurring in the implementation of
a SYCL kernel invocation function such as `sycl::handler::single_task()`.
  template<typename KernelNameType, typename KernelType>
  [[clang::sycl_kernel_entry_point(KernelNameType)]]
  void kernel_entry_point(KernelType kerne) {
    kernel();
  }

The body of the above function specifies the parameters and body of the
generated offload kernel entry point. Clearly, a call to the above function
by a SYCL kernel invocation function is not intended to execute the body
as written. Previously, code generation emitted an empty function body so
that calls to the function had no effect other than to trigger the generation
of the offload kernel entry point. The function body is therefore available
to hook for SYCL library support and is now substituted with a call to a
(SYCL library provided) function template named `sycl_enqueue_kernel_launch()`
with the kernel name type passed as the first template argument, the
symbol name of the offload kernel entry point passed as a string literal for
the first function argument, and the (possibly decomposed) parameters passed
as the remaining explicit function arguments. Given a call like this:
  kernel_entry_point<struct KN>([]{})
the body of the instantiated `kernel_entry_point()` specialization would be
substituted as follows with "kernel-symbol-name" substituted for the
generated symbol name and `kernel` forwarded (This assumes no kernel
argument decomposition; if decomposition was required, `kernel` would be
replaced with its corresponding decomposed arguments).
  sycl_enqueue_kernel_launch<KN>("kernel-symbol-name", kernel)

Name lookup and overload resolution for the `sycl_enqueue_kernel_launch()`
function is performed at the point of definition of the
`sycl_kernel_entry_point` attributed function (or the point of instantiation
for an instantiated function template specialization). If overload
resolution fails, the program is ill-formed.

Implementation of the `sycl_enqueue_kernel_launch()` function might require
additional information provided by the SYCL library. This is facilitated by
removing the previous prohibition against use of the `sycl_kernel_entry_point`
attribute with a non-static member function. If the `sycl_kernel_entry_point`
attributed function is a non-static member function, then overload resolution
for the `sycl_enqueue_kernel_launch()` function template may select a
non-static member function in which case, `this` will be implicitly passed
as the implicit object argument.

If a `sycl_kernel_entry_point` attributed function is a non-static member
function, use of `this` in a potentially evaluated expression is prohibited
in the definition (since `this` is not a kernel argument and will not be
available within the generated offload kernel entry point function).

Support for kernel argument decomposition and reconstruction is not yet
implemented.

Author: tahonermann

clang/include/clang/AST/ASTNodeTraverser.h

@@ -836,8 +836,10 @@ class ASTNodeTraverser
 
   void VisitSYCLKernelCallStmt(const SYCLKernelCallStmt *Node) {
     Visit(Node->getOriginalStmt());
-    if (Traversal != TK_IgnoreUnlessSpelledInSource)
+    if (Traversal != TK_IgnoreUnlessSpelledInSource) {
+      Visit(Node->getKernelLaunchStmt());
       Visit(Node->getOutlinedFunctionDecl());
+    }
   }
 
   void VisitOMPExecutableDirective(const OMPExecutableDirective *Node) {

clang/include/clang/AST/RecursiveASTVisitor.h

@@ -3027,6 +3027,7 @@ DEF_TRAVERSE_STMT(CapturedStmt, { TRY_TO(TraverseDecl(S->getCapturedDecl())); })
 DEF_TRAVERSE_STMT(SYCLKernelCallStmt, {
   if (getDerived().shouldVisitImplicitCode()) {
     TRY_TO(TraverseStmt(S->getOriginalStmt()));
+    TRY_TO(TraverseStmt(S->getKernelLaunchStmt()));
     TRY_TO(TraverseDecl(S->getOutlinedFunctionDecl()));
     ShouldVisitChildren = false;
   }

clang/include/clang/AST/StmtSYCL.h

@@ -28,35 +28,45 @@ namespace clang {
 /// of such a function specifies the statements to be executed on a SYCL device
 /// to invoke a SYCL kernel with a particular set of kernel arguments. The
 /// SYCLKernelCallStmt associates an original statement (the compound statement
-/// that is the function body) with an OutlinedFunctionDecl that holds the
-/// kernel parameters and the transformed body. During code generation, the
-/// OutlinedFunctionDecl is used to emit an offload kernel entry point suitable
-/// for invocation from a SYCL library implementation. If executed, the
-/// SYCLKernelCallStmt behaves as a no-op; no code generation is performed for
-/// it.
+/// that is the function body) with a kernel launch statement to execute on a
+/// SYCL host and an OutlinedFunctionDecl that holds the kernel parameters and
+/// the transformed body to execute on a SYCL device. During code generation,
+/// the OutlinedFunctionDecl is used to emit an offload kernel entry point
+/// suitable for invocation from a SYCL library implementation.
 class SYCLKernelCallStmt : public Stmt {
   friend class ASTStmtReader;
   friend class ASTStmtWriter;
 
 private:
   Stmt *OriginalStmt = nullptr;
+  Stmt *KernelLaunchStmt = nullptr;
   OutlinedFunctionDecl *OFDecl = nullptr;
 
 public:
   /// Construct a SYCL kernel call statement.
-  SYCLKernelCallStmt(CompoundStmt *CS, OutlinedFunctionDecl *OFD)
-      : Stmt(SYCLKernelCallStmtClass), OriginalStmt(CS), OFDecl(OFD) {}
+  SYCLKernelCallStmt(CompoundStmt *CS, Stmt *S, OutlinedFunctionDecl *OFD)
+      : Stmt(SYCLKernelCallStmtClass), OriginalStmt(CS), KernelLaunchStmt(S),
+        OFDecl(OFD) {}
 
   /// Construct an empty SYCL kernel call statement.
   SYCLKernelCallStmt(EmptyShell Empty) : Stmt(SYCLKernelCallStmtClass, Empty) {}
 
-  /// Retrieve the model statement.
+  /// Retrieve the original statement.
   CompoundStmt *getOriginalStmt() { return cast<CompoundStmt>(OriginalStmt); }
   const CompoundStmt *getOriginalStmt() const {
     return cast<CompoundStmt>(OriginalStmt);
   }
+
+  /// Set the original statement.
   void setOriginalStmt(CompoundStmt *CS) { OriginalStmt = CS; }
 
+  /// Retrieve the kernel launch statement.
+  Stmt *getKernelLaunchStmt() { return KernelLaunchStmt; }
+  const Stmt *getKernelLaunchStmt() const { return KernelLaunchStmt; }
+
+  /// Set the kernel launch statement.
+  void setKernelLaunchStmt(Stmt *S) { KernelLaunchStmt = S; }
+
   /// Retrieve the outlined function declaration.
   OutlinedFunctionDecl *getOutlinedFunctionDecl() { return OFDecl; }
   const OutlinedFunctionDecl *getOutlinedFunctionDecl() const { return OFDecl; }

clang/include/clang/Basic/AttrDocs.td

@@ -532,13 +532,13 @@ The following examples demonstrate the use of this attribute:
 def SYCLKernelEntryPointDocs : Documentation {
   let Category = DocCatFunction;
   let Content = [{
-The ``sycl_kernel_entry_point`` attribute facilitates the generation of an
-offload kernel entry point, sometimes called a SYCL kernel caller function,
-suitable for invoking a SYCL kernel on an offload device. The attribute is
-intended for use in the implementation of SYCL kernel invocation functions
-like the ``single_task`` and ``parallel_for`` member functions of the
-``sycl::handler`` class specified in section 4.9.4, "Command group ``handler``
-class", of the SYCL 2020 specification.
+The ``sycl_kernel_entry_point`` attribute facilitates the launch of a SYCL
+kernel and the generation of an offload kernel entry point, sometimes called
+a SYCL kernel caller function, suitable for invoking a SYCL kernel on an
+offload device. The attribute is intended for use in the implementation of
+SYCL kernel invocation functions like the ``single_task`` and ``parallel_for``
+member functions of the ``sycl::handler`` class specified in section 4.9.4,
+"Command group ``handler`` class", of the SYCL 2020 specification.
 
 The attribute requires a single type argument that specifies a class type that
 meets the requirements for a SYCL kernel name as described in section 5.2,
@@ -550,7 +550,7 @@ The attribute only appertains to functions and only those that meet the
 following requirements.
 
 * Has a non-deduced ``void`` return type.
-* Is not a non-static member function, constructor, or destructor.
+* Is not a constructor or destructor.
 * Is not a C variadic function.
 * Is not a coroutine.
 * Is not defined as deleted or as defaulted.
@@ -565,39 +565,43 @@ follows.
 
   namespace sycl {
   class handler {
+    template<typename KernelNameType, typename... Ts>
+    void sycl_enqueue_kernel_launch(const char *KernelName, Ts...) {
+      // Call functions appropriate for the desired offload backend
+      // (OpenCL, CUDA, HIP, Level Zero, etc...) to enqueue kernel invocation.
+    }
+
     template<typename KernelNameType, typename KernelType>
     [[ clang::sycl_kernel_entry_point(KernelNameType) ]]
-    static void kernel_entry_point(KernelType kernel) {
-      kernel();
+    void kernel_entry_point(KernelType Kernel) {
+      Kernel();
     }
 
   public:
     template<typename KernelNameType, typename KernelType>
-    void single_task(KernelType kernel) {
-      // Call kernel_entry_point() to trigger generation of an offload
-      // kernel entry point.
-      kernel_entry_point<KernelNameType>(kernel);
-      // Call functions appropriate for the desired offload backend
-      // (OpenCL, CUDA, HIP, Level Zero, etc...).
+    void single_task(KernelType Kernel) {
+      // Call kernel_entry_point() to launch the kernel and to trigger
+      // generation of an offload kernel entry point.
+      kernel_entry_point<KernelNameType>(Kernel);
     }
   };
   } // namespace sycl
 
-A SYCL kernel is a callable object of class type that is constructed on a host,
-often via a lambda expression, and then passed to a SYCL kernel invocation
-function to be executed on an offload device. A SYCL kernel invocation function
-is responsible for copying the provided SYCL kernel object to an offload
-device and initiating a call to it. The SYCL kernel object and its data members
-constitute the parameters of an offload kernel.
-
-A SYCL kernel type is required to satisfy the device copyability requirements
-specified in section 3.13.1, "Device copyable", of the SYCL 2020 specification.
-Additionally, any data members of the kernel object type are required to satisfy
-section 4.12.4, "Rules for parameter passing to kernels". For most types, these
-rules require that the type is trivially copyable.  However, the SYCL
-specification mandates that certain special SYCL types, such as
-``sycl::accessor`` and ``sycl::stream`` be device copyable even if they are not
-trivially copyable. These types require special handling because they cannot
+A SYCL kernel object is a callable object of class type that is constructed on
+a host, often via a lambda expression, and then passed to a SYCL kernel
+invocation function to be executed on an offload device. A SYCL kernel
+invocation function is responsible for copying the provided SYCL kernel object
+to an offload device and initiating a call to it. The SYCL kernel object and
+its data members constitute the parameters of an offload kernel.
+
+A SYCL kernel object type is required to satisfy the device copyability
+requirements specified in section 3.13.1, "Device copyable", of the SYCL 2020
+specification. Additionally, any data members of the kernel object type are
+required to satisfy section 4.12.4, "Rules for parameter passing to kernels".
+For most types, these rules require that the type is trivially copyable.
+However, the SYCL specification mandates that certain special SYCL types, such
+as ``sycl::accessor`` and ``sycl::stream`` be device copyable even if they are
+not trivially copyable. These types require special handling because they cannot
 be copied to device memory as if by ``memcpy()``. Additionally, some offload
 backends, OpenCL for example, require objects of some of these types to be
 passed as individual arguments to the offload kernel.
@@ -612,7 +616,7 @@ like OpenCL):
 
 #. Identifying the offload kernel entry point to be used for the SYCL kernel.
 
-#. Deconstructing the SYCL kernel object, if necessary, to produce the set of
+#. Decomposing the SYCL kernel object, if necessary, to produce the set of
    offload kernel arguments required by the offload kernel entry point.
 
 #. Copying the offload kernel arguments to device memory.
@@ -621,17 +625,23 @@ like OpenCL):
 
 The offload kernel entry point for a SYCL kernel performs the following tasks:
 
-#. Reconstituting the SYCL kernel object, if necessary, using the offload
+#. Reconstructing the SYCL kernel object, if necessary, using the offload
    kernel parameters.
 
-#. Calling the ``operator()`` member function of the (reconstituted) SYCL kernel
+#. Calling the ``operator()`` member function of the (reconstructed) SYCL kernel
    object.
 
-The ``sycl_kernel_entry_point`` attribute automates generation of an offload
-kernel entry point that performs those latter tasks. The parameters and body of
-a function declared with the ``sycl_kernel_entry_point`` attribute specify a
-pattern from which the parameters and body of the entry point function are
-derived. Consider the following call to a SYCL kernel invocation function.
+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
+with the kernel name type, kernel symbol name, and (decomposed) kernel arguments
+passed as template or function arguments.
+
+A function declared with the ``sycl_kernel_entry_point`` attribute specifies
+the parameters and body of the offload entry point function. Consider the
+following call to the ``single_task()`` SYCL kernel invocation function assuming
+an implementation similar to the one shown above.
 
 .. code-block:: c++
 
@@ -645,31 +655,33 @@ derived. Consider the following call to a SYCL kernel invocation function.
 The SYCL kernel object is the result of the lambda expression. It has two
 data members corresponding to the captures of ``sout`` and ``s``. Since one
 of these data members corresponds to a special SYCL type that must be passed
-individually as an offload kernel parameter, it is necessary to decompose the
-SYCL kernel object into its constituent parts; the offload kernel will have
-two kernel parameters. Given a SYCL implementation that uses a
-``sycl_kernel_entry_point`` attributed function like the one shown above, an
+individually as an offload kernel argument, it is necessary to decompose the
+SYCL kernel object into its constituent parts and pass them individually. An
 offload kernel entry point function will be generated that looks approximately
 as follows.
 
 .. code-block:: c++
 
   void sycl-kernel-caller-for-KN(sycl::stream sout, S s) {
-    kernel-type kernel = { sout, s );
-    kernel();
+    kernel-type Kernel = { sout, s );
+    Kernel();
   }
 
 There are a few items worthy of note:
 
 #. The name of the generated function incorporates the SYCL kernel name,
    ``KN``, that was passed as the ``KernelNameType`` template parameter to
-   ``kernel_entry_point()`` and provided as the argument to the
+   ``single_task()`` and eventually provided as the argument to the
    ``sycl_kernel_entry_point`` attribute. There is a one-to-one correspondence
    between SYCL kernel names and offload kernel entry points.
 
+#. The parameters and the call to ``Kernel()`` correspond to the definition of
+   ``kernel_entry_point()`` called by ``single_task()`` with the SYCL kernel
+   object argument decomposed and reconstructed.
+
 #. The SYCL kernel is a lambda closure type and therefore has no name;
    ``kernel-type`` is substituted above and corresponds to the ``KernelType``
-   template parameter deduced in the call to ``kernel_entry_point()``.
+   template parameter deduced in the call to ``single_task()``.
    Lambda types cannot be declared and initialized using the aggregate
    initialization syntax used above, but the intended behavior should be clear.
 
@@ -683,24 +695,55 @@ There are a few items worthy of note:
    or more parameters depending on how the SYCL library implementation defines
    these types.
 
-#. The call to ``kernel_entry_point()`` has no effect other than to trigger
-   emission of the entry point function. The statments that make up the body
-   of the function are not executed when the function is called; they are
-   only used in the generation of the entry point function.
+The call to ``kernel_entry_point()`` by ``single_task()`` is effectively
+replaced with synthesized code that looks approximately as follows.
+
+.. code-block:: c++
+
+  sycl::stream sout = Kernel.sout;
+  S s = Kernel.s;
+  sycl_enqueue_kernel_launch<KN>("kernel-symbol-name", sout, s);
+
+There are a few items worthy of note:
+
+#. The SYCL kernel object is a lambda closure type and its captures do not
+   have formal names and cannot be accessed using the member access syntax used
+   above, but the intended behavior should be clear.
+
+#. ``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
+   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
+   as if by ``std::forward<X>(x)``.
+
+#. The ``sycl_enqueue_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
+   device memory, presumably via an offload backend such as OpenCL.
 
 It is not necessary for a function declared with the ``sycl_kernel_entry_point``
 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.
+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.
 
 Functions declared with the ``sycl_kernel_entry_point`` attribute are not
 limited to the simple example shown above. They may have additional template
 parameters, declare additional function parameters, and have complex control
-flow in the function body. Function parameter decomposition and reconstitution
+flow in the function body. Function parameter decomposition and reconstruction
 is performed for all function parameters. The function must abide by the
 language feature restrictions described in section 5.4, "Language restrictions
-for device functions" in the SYCL 2020 specification.
+for device functions" in the SYCL 2020 specification. If the function is a
+non-static member function, ``this`` shall not be used in a potentially
+evaluated expression.
   }];
 }
 

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -13062,9 +13062,10 @@ def warn_sycl_external_missing_on_first_decl : Warning<
 // SYCL kernel entry point diagnostics
 def err_sycl_entry_point_invalid : Error<
   "the %0 attribute cannot be applied to a"
-  " %select{non-static member function|variadic function|deleted function|"
-           "defaulted function|constexpr function|consteval function|"
-           "function declared with the 'noreturn' attribute|coroutine|"
+  " %select{variadic function|deleted function|defaulted function|"
+           "constructor|destructor|coroutine|"
+           "constexpr function|consteval function|"
+           "function declared with the 'noreturn' attribute|"
            "function defined with a function try block}1">;
 def err_sycl_entry_point_invalid_redeclaration : Error<
   "the %0 kernel name argument does not match prior"

clang/lib/AST/ASTContext.cpp

@@ -14975,6 +14975,10 @@ static SYCLKernelInfo BuildSYCLKernelInfo(ASTContext &Context,
   MC->mangleCanonicalTypeName(KernelNameType, Out);
   std::string KernelName = Out.str();
 
+  // FIXME: Diagnose kernel names that are not representable in the ordinary
+  // literal encoding. This is not necessarily the right place to add such
+  // a diagnostic.
+
   return {KernelNameType, FD, KernelName};
 }
 

clang/lib/AST/StmtPrinter.cpp

@@ -595,7 +595,7 @@ void StmtPrinter::VisitCapturedStmt(CapturedStmt *Node) {
 }
 
 void StmtPrinter::VisitSYCLKernelCallStmt(SYCLKernelCallStmt *Node) {
-  PrintStmt(Node->getOutlinedFunctionDecl()->getBody());
+  PrintStmt(Node->getOriginalStmt());
 }
 
 void StmtPrinter::VisitObjCAtTryStmt(ObjCAtTryStmt *Node) {

clang/lib/CodeGen/CGStmt.cpp

@@ -19,6 +19,7 @@
 #include "clang/AST/Attr.h"
 #include "clang/AST/Expr.h"
 #include "clang/AST/Stmt.h"
+#include "clang/AST/StmtSYCL.h"
 #include "clang/AST/StmtVisitor.h"
 #include "clang/Basic/Builtins.h"
 #include "clang/Basic/DiagnosticSema.h"
@@ -543,21 +544,7 @@ bool CodeGenFunction::EmitSimpleStmt(const Stmt *S,
     EmitSEHLeaveStmt(cast<SEHLeaveStmt>(*S));
     break;
   case Stmt::SYCLKernelCallStmtClass:
-    // SYCL kernel call statements are generated as wrappers around the body
-    // of functions declared with the sycl_kernel_entry_point attribute. Such
-    // functions are used to specify how a SYCL kernel (a function object) is
-    // to be invoked; the SYCL kernel call statement contains a transformed
-    // variation of the function body and is used to generate a SYCL kernel
-    // caller function; a function that serves as the device side entry point
-    // used to execute the SYCL kernel. The sycl_kernel_entry_point attributed
-    // function is invoked by host code in order to trigger emission of the
-    // device side SYCL kernel caller function and to generate metadata needed
-    // by SYCL run-time library implementations; the function is otherwise
-    // intended to have no effect. As such, the function body is not evaluated
-    // as part of the invocation during host compilation (and the function
-    // should not be called or emitted during device compilation); the SYCL
-    // kernel call statement is thus handled as a null statement for the
-    // purpose of code generation.
+    EmitSYCLKernelCallStmt(cast<SYCLKernelCallStmt>(*S));
     break;
   }
   return true;

clang/lib/CodeGen/CodeGenFunction.h

@@ -3643,6 +3643,8 @@ class CodeGenFunction : public CodeGenTypeCache {
   LValue EmitCoyieldLValue(const CoyieldExpr *E);
   RValue EmitCoroutineIntrinsic(const CallExpr *E, unsigned int IID);
 
+  void EmitSYCLKernelCallStmt(const SYCLKernelCallStmt &S);
+
   void EnterCXXTryStmt(const CXXTryStmt &S, bool IsFnTryBlock = false);
   void ExitCXXTryStmt(const CXXTryStmt &S, bool IsFnTryBlock = false);