https://github.com/tahonermann updated
https://github.com/llvm/llvm-project/pull/152403
>From 5b42f6b9470a96a79b3474fcf384fa950449e095 Mon Sep 17 00:00:00 2001
From: Tom Honermann <tom.honerm...@intel.com>
Date: Wed, 6 Aug 2025 19:26:54 -0700
Subject: [PATCH] [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.
---
clang/include/clang/AST/ASTNodeTraverser.h | 4 +-
clang/include/clang/AST/RecursiveASTVisitor.h | 1 +
clang/include/clang/AST/StmtSYCL.h | 28 ++--
clang/include/clang/Basic/AttrDocs.td | 151 +++++++++++-------
.../clang/Basic/DiagnosticSemaKinds.td | 7 +-
clang/lib/AST/ASTContext.cpp | 4 +
clang/lib/AST/StmtPrinter.cpp | 2 +-
clang/lib/CodeGen/CGStmt.cpp | 17 +-
clang/lib/CodeGen/CodeGenFunction.h | 2 +
clang/lib/CodeGen/CodeGenSYCL.cpp | 15 ++
clang/lib/Sema/SemaDecl.cpp | 8 +-
clang/lib/Sema/SemaExceptionSpec.cpp | 11 +-
clang/lib/Sema/SemaSYCL.cpp | 151 ++++++++++++++----
clang/lib/Serialization/ASTReaderStmt.cpp | 1 +
clang/lib/Serialization/ASTWriterStmt.cpp | 1 +
clang/test/AST/ast-print-sycl-kernel-call.cpp | 22 +++
.../ast-dump-sycl-kernel-call-stmt.cpp | 51 +++++-
.../CodeGenSYCL/kernel-caller-entry-point.cpp | 13 +-
...-kernel-entry-point-attr-appertainment.cpp | 29 ++--
.../sycl-kernel-entry-point-attr-grammar.cpp | 2 +
...cl-kernel-entry-point-attr-kernel-name.cpp | 2 +
.../sycl-kernel-entry-point-attr-sfinae.cpp | 2 +
22 files changed, 385 insertions(+), 139 deletions(-)
create mode 100644 clang/test/AST/ast-print-sycl-kernel-call.cpp
diff --git a/clang/include/clang/AST/ASTNodeTraverser.h
b/clang/include/clang/AST/ASTNodeTraverser.h
index 8ebabb2bde10d..ab12b659e34ba 100644
--- a/clang/include/clang/AST/ASTNodeTraverser.h
+++ b/clang/include/clang/AST/ASTNodeTraverser.h
@@ -848,8 +848,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) {
diff --git a/clang/include/clang/AST/RecursiveASTVisitor.h
b/clang/include/clang/AST/RecursiveASTVisitor.h
index 62991d986e675..b9fa412a49685 100644
--- a/clang/include/clang/AST/RecursiveASTVisitor.h
+++ b/clang/include/clang/AST/RecursiveASTVisitor.h
@@ -2942,6 +2942,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;
}
diff --git a/clang/include/clang/AST/StmtSYCL.h
b/clang/include/clang/AST/StmtSYCL.h
index 28ace12d7916b..70d8137992110 100644
--- a/clang/include/clang/AST/StmtSYCL.h
+++ b/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;
}
diff --git a/clang/include/clang/Basic/AttrDocs.td
b/clang/include/clang/Basic/AttrDocs.td
index 2b095ab975202..cf8241ab79ae6 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -479,13 +479,13 @@ The SYCL kernel in the previous code sample meets these
expectations.
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,
@@ -497,7 +497,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.
@@ -512,39 +512,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.
@@ -559,7 +563,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.
@@ -568,17 +572,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++
@@ -592,31 +602,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.
@@ -630,24 +642,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.
}];
}
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td
b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index cf23594201143..58f2d9d5bd250 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12945,9 +12945,10 @@ def err_sycl_special_type_num_init_method : Error<
// 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"
diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp
index 3a16111dd5f7d..4636a318a278a 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/clang/lib/AST/ASTContext.cpp
@@ -15034,6 +15034,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};
}
diff --git a/clang/lib/AST/StmtPrinter.cpp b/clang/lib/AST/StmtPrinter.cpp
index 6ba5ec89964a9..363a9cb2c7815 100644
--- a/clang/lib/AST/StmtPrinter.cpp
+++ b/clang/lib/AST/StmtPrinter.cpp
@@ -589,7 +589,7 @@ void StmtPrinter::VisitCapturedStmt(CapturedStmt *Node) {
}
void StmtPrinter::VisitSYCLKernelCallStmt(SYCLKernelCallStmt *Node) {
- PrintStmt(Node->getOutlinedFunctionDecl()->getBody());
+ PrintStmt(Node->getOriginalStmt());
}
void StmtPrinter::VisitObjCAtTryStmt(ObjCAtTryStmt *Node) {
diff --git a/clang/lib/CodeGen/CGStmt.cpp b/clang/lib/CodeGen/CGStmt.cpp
index 1a8c6f015bda1..da05d68479e0f 100644
--- a/clang/lib/CodeGen/CGStmt.cpp
+++ b/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"
@@ -540,21 +541,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;
diff --git a/clang/lib/CodeGen/CodeGenFunction.h
b/clang/lib/CodeGen/CodeGenFunction.h
index 6c32c98cec011..60e67325fcb52 100644
--- a/clang/lib/CodeGen/CodeGenFunction.h
+++ b/clang/lib/CodeGen/CodeGenFunction.h
@@ -3626,6 +3626,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);
diff --git a/clang/lib/CodeGen/CodeGenSYCL.cpp
b/clang/lib/CodeGen/CodeGenSYCL.cpp
index b9a96fe8ab838..29a9659d0d14d 100644
--- a/clang/lib/CodeGen/CodeGenSYCL.cpp
+++ b/clang/lib/CodeGen/CodeGenSYCL.cpp
@@ -17,6 +17,21 @@
using namespace clang;
using namespace CodeGen;
+void CodeGenFunction::EmitSYCLKernelCallStmt(const SYCLKernelCallStmt &S) {
+ if (getLangOpts().SYCLIsDevice) {
+ // A sycl_kernel_entry_point attributed function is unlikely to be emitted
+ // during device compilation, but might be if it is ODR-used from device
+ // code that is emitted. In these cases, the function is emitted with an
+ // empty body; the original body is emitted in the offload kernel entry
+ // point and the synthesized kernel launch code is only relevant for host
+ // compilation.
+ return;
+ }
+
+ assert(getLangOpts().SYCLIsHost);
+ EmitStmt(S.getKernelLaunchStmt());
+}
+
static void SetSYCLKernelAttributes(llvm::Function *Fn, CodeGenFunction &CGF) {
// SYCL 2020 device language restrictions require forward progress and
// disallow recursion.
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index b5eb825eb52cc..50a6278688948 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -16255,19 +16255,19 @@ Decl *Sema::ActOnFinishFunctionBody(Decl *dcl, Stmt
*Body,
FD->getAttr<SYCLKernelEntryPointAttr>();
if (FD->isDefaulted()) {
Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid)
- << SKEPAttr << /*defaulted function*/ 3;
+ << SKEPAttr << /*defaulted function*/ 2;
SKEPAttr->setInvalidAttr();
} else if (FD->isDeleted()) {
Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid)
- << SKEPAttr << /*deleted function*/ 2;
+ << SKEPAttr << /*deleted function*/ 1;
SKEPAttr->setInvalidAttr();
} else if (FSI->isCoroutine()) {
Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid)
- << SKEPAttr << /*coroutine*/ 7;
+ << SKEPAttr << /*coroutine*/ 5;
SKEPAttr->setInvalidAttr();
} else if (Body && isa<CXXTryStmt>(Body)) {
Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid)
- << SKEPAttr << /*function defined with a function try block*/ 8;
+ << SKEPAttr << /*function defined with a function try block*/ 9;
SKEPAttr->setInvalidAttr();
}
diff --git a/clang/lib/Sema/SemaExceptionSpec.cpp
b/clang/lib/Sema/SemaExceptionSpec.cpp
index 0a6cea8869c14..91e4ce2790d0c 100644
--- a/clang/lib/Sema/SemaExceptionSpec.cpp
+++ b/clang/lib/Sema/SemaExceptionSpec.cpp
@@ -15,6 +15,7 @@
#include "clang/AST/Expr.h"
#include "clang/AST/ExprCXX.h"
#include "clang/AST/StmtObjC.h"
+#include "clang/AST/StmtSYCL.h"
#include "clang/AST/TypeLoc.h"
#include "clang/Basic/Diagnostic.h"
#include "clang/Basic/SourceManager.h"
@@ -1250,6 +1251,15 @@ CanThrowResult Sema::canThrow(const Stmt *S) {
return CT;
}
+ case Stmt::SYCLKernelCallStmtClass: {
+ auto *SKCS = cast<SYCLKernelCallStmt>(S);
+ if (getLangOpts().SYCLIsDevice)
+ return canSubStmtsThrow(*this,
+ SKCS->getOutlinedFunctionDecl()->getBody());
+ assert(getLangOpts().SYCLIsHost);
+ return canSubStmtsThrow(*this, SKCS->getKernelLaunchStmt());
+ }
+
// ObjC message sends are like function calls, but never have exception
// specs.
case Expr::ObjCMessageExprClass:
@@ -1430,7 +1440,6 @@ CanThrowResult Sema::canThrow(const Stmt *S) {
case Stmt::AttributedStmtClass:
case Stmt::BreakStmtClass:
case Stmt::CapturedStmtClass:
- case Stmt::SYCLKernelCallStmtClass:
case Stmt::CaseStmtClass:
case Stmt::CompoundStmtClass:
case Stmt::ContinueStmtClass:
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 4683c81bd1c60..eec087b64ac47 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -298,43 +298,46 @@ void
SemaSYCL::CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD) {
}
}
- if (const auto *MD = dyn_cast<CXXMethodDecl>(FD)) {
- if (!MD->isStatic()) {
- Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid)
- << SKEPAttr << /*non-static member function*/ 0;
- SKEPAttr->setInvalidAttr();
- }
+ if (isa<CXXConstructorDecl>(FD)) {
+ Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid)
+ << SKEPAttr << /*constructor*/ 3;
+ SKEPAttr->setInvalidAttr();
+ }
+ if (isa<CXXDestructorDecl>(FD)) {
+ Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid)
+ << SKEPAttr << /*destructor*/ 4;
+ SKEPAttr->setInvalidAttr();
}
if (FD->isVariadic()) {
Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid)
- << SKEPAttr << /*variadic function*/ 1;
+ << SKEPAttr << /*variadic function*/ 0;
SKEPAttr->setInvalidAttr();
}
if (FD->isDefaulted()) {
Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid)
- << SKEPAttr << /*defaulted function*/ 3;
+ << SKEPAttr << /*defaulted function*/ 2;
SKEPAttr->setInvalidAttr();
} else if (FD->isDeleted()) {
Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid)
- << SKEPAttr << /*deleted function*/ 2;
+ << SKEPAttr << /*deleted function*/ 1;
SKEPAttr->setInvalidAttr();
}
if (FD->isConsteval()) {
Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid)
- << SKEPAttr << /*consteval function*/ 5;
+ << SKEPAttr << /*consteval function*/ 7;
SKEPAttr->setInvalidAttr();
} else if (FD->isConstexpr()) {
Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid)
- << SKEPAttr << /*constexpr function*/ 4;
+ << SKEPAttr << /*constexpr function*/ 6;
SKEPAttr->setInvalidAttr();
}
if (FD->isNoReturn()) {
Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid)
- << SKEPAttr << /*function declared with the 'noreturn' attribute*/ 6;
+ << SKEPAttr << /*function declared with the 'noreturn' attribute*/ 8;
SKEPAttr->setInvalidAttr();
}
@@ -372,6 +375,67 @@ void
SemaSYCL::CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD) {
namespace {
+CompoundStmt *BuildSYCLKernelLaunchStmt(Sema &SemaRef, FunctionDecl *FD,
+ const std::string &KernelName) {
+ ASTContext &Ctx = SemaRef.getASTContext();
+ SmallVector<Stmt *> Stmts;
+
+ // Prepare a string literal that contains the kernel name in the ordinary
+ // literal encoding.
+ // FIXME: transcode the contents of KernelName from UTF-8 to the
+ // ordinary literal encoding.
+ QualType KernelNameCharTy = Ctx.CharTy.withConst();
+ llvm::APInt KernelNameSize(Ctx.getTypeSize(Ctx.getSizeType()),
+ KernelName.size() + 1);
+ QualType KernelNameArrayTy = Ctx.getConstantArrayType(
+ KernelNameCharTy, KernelNameSize, nullptr, ArraySizeModifier::Normal, 0);
+ StringLiteral *KernelNameExpr = StringLiteral::Create(
+ Ctx, KernelName, StringLiteralKind::Ordinary,
+ /*Pascal*/ false, KernelNameArrayTy, SourceLocation());
+
+ // FIXME: An extern variable declaration with assignment to the kernel
+ // name expression is added to Stmts as a temporary measure to see results.
+ // reflected in tests. The kernel name expression will need to be passed as
+ // the first function argument in a call to sycl_enqueue_kernel_launch.
+ QualType ExternVarType = Ctx.getPointerType(Ctx.CharTy.withConst());
+ const IdentifierInfo *ExternVarName =
+ SemaRef.getPreprocessor().getIdentifierInfo("kernel_name");
+ VarDecl *ExternVarDecl = VarDecl::Create(
+ Ctx, FD, SourceLocation(), SourceLocation(), ExternVarName,
ExternVarType,
+ /*TInfo*/ nullptr, SC_Extern);
+ DeclStmt *ExternVarDeclStmt = new (Ctx)
+ DeclStmt(DeclGroupRef(ExternVarDecl), SourceLocation(),
SourceLocation());
+ Stmts.push_back(ExternVarDeclStmt);
+ DeclRefExpr *ExternVarDeclRef = new (Ctx) DeclRefExpr(
+ Ctx, ExternVarDecl, /*RefersToEnclosingVariableOrCapture*/ false,
+ ExternVarType, VK_LValue, SourceLocation());
+ ImplicitCastExpr *KernelNameArrayDecayExpr = new (Ctx) ImplicitCastExpr(
+ ImplicitCastExpr::OnStack, ExternVarType, CK_ArrayToPointerDecay,
+ KernelNameExpr, VK_PRValue, FPOptionsOverride());
+ BinaryOperator *AssignmentExpr = BinaryOperator::Create(
+ Ctx, ExternVarDeclRef, KernelNameArrayDecayExpr, BO_Assign,
ExternVarType,
+ VK_LValue, OK_Ordinary, SourceLocation(), FPOptionsOverride());
+ Stmts.push_back(AssignmentExpr);
+
+ // Perform overload resolution for a call to an accessible (member) function
+ // template named 'sycl_enqueue_kernel_launch' from within the definition of
+ // FD where:
+ // - The kernel name type is passed as the first template argument.
+ // - Any remaining template parameters are deduced from the function
arguments
+ // or assigned by default template arguments.
+ // - 'this' is passed as the implicit function argument if 'FD' is a
+ // non-static member function.
+ // - The name of the kernel, expressed as a string literal, is passed as the
+ // first function argument.
+ // - The parameters of FD are forwarded as-if by 'std::forward()' as the
+ // remaining explicit function arguments.
+ // - Any remaining function arguments are initialized by default arguments.
+ CompoundStmt *LaunchStmt = CompoundStmt::Create(
+ Ctx, Stmts, FPOptionsOverride(), SourceLocation(), SourceLocation());
+
+ return LaunchStmt;
+}
+
// The body of a function declared with the [[sycl_kernel_entry_point]]
// attribute is cloned and transformed to substitute references to the original
// function parameters with references to replacement variables that stand in
@@ -415,6 +479,36 @@ class OutlinedFunctionDeclBodyInstantiator
ParmDeclMap &MapRef;
};
+OutlinedFunctionDecl *BuildSYCLKernelEntryPointOutline(Sema &SemaRef,
+ FunctionDecl *FD,
+ CompoundStmt *Body) {
+ using ParmDeclMap = OutlinedFunctionDeclBodyInstantiator::ParmDeclMap;
+ ParmDeclMap ParmMap;
+
+ OutlinedFunctionDecl *OFD = OutlinedFunctionDecl::Create(
+ SemaRef.getASTContext(), FD, FD->getNumParams());
+ unsigned i = 0;
+ for (ParmVarDecl *PVD : FD->parameters()) {
+ ImplicitParamDecl *IPD = ImplicitParamDecl::Create(
+ SemaRef.getASTContext(), OFD, SourceLocation(), PVD->getIdentifier(),
+ PVD->getType(), ImplicitParamKind::Other);
+ OFD->setParam(i, IPD);
+ ParmMap[PVD] = IPD;
+ ++i;
+ }
+
+ // FIXME: Diagnose (implicit or explicit) use of CXXThisExpr in potentially
+ // evaluated contexts in the function body. This is not necessarily the
+ // right place to add such a diagnostic.
+
+ OutlinedFunctionDeclBodyInstantiator OFDBodyInstantiator(SemaRef, ParmMap);
+ Stmt *OFDBody = OFDBodyInstantiator.TransformStmt(Body).get();
+ OFD->setBody(OFDBody);
+ OFD->setNothrow();
+
+ return OFD;
+}
+
} // unnamed namespace
StmtResult SemaSYCL::BuildSYCLKernelCallStmt(FunctionDecl *FD,
@@ -423,6 +517,11 @@ StmtResult SemaSYCL::BuildSYCLKernelCallStmt(FunctionDecl
*FD,
assert(!FD->isTemplated());
assert(FD->hasPrototype());
+ // The current context must be the function definition context to ensure
+ // that name lookup and parameter and local variable creation are performed
+ // within the correct scope.
+ assert(SemaRef.CurContext == FD);
+
const auto *SKEPAttr = FD->getAttr<SYCLKernelEntryPointAttr>();
assert(SKEPAttr && "Missing sycl_kernel_entry_point attribute");
assert(!SKEPAttr->isInvalidAttr() &&
@@ -434,29 +533,19 @@ StmtResult SemaSYCL::BuildSYCLKernelCallStmt(FunctionDecl
*FD,
getASTContext().getSYCLKernelInfo(SKEPAttr->getKernelName());
assert(declaresSameEntity(SKI.getKernelEntryPointDecl(), FD) &&
"SYCL kernel name conflict");
- (void)SKI;
- using ParmDeclMap = OutlinedFunctionDeclBodyInstantiator::ParmDeclMap;
- ParmDeclMap ParmMap;
+ // Build the kernel launch statement.
+ Stmt *LaunchStmt =
+ BuildSYCLKernelLaunchStmt(SemaRef, FD, SKI.GetKernelName());
+ assert(LaunchStmt);
- assert(SemaRef.CurContext == FD);
+ // Build the outline of the synthesized device entry point function.
OutlinedFunctionDecl *OFD =
- OutlinedFunctionDecl::Create(getASTContext(), FD, FD->getNumParams());
- unsigned i = 0;
- for (ParmVarDecl *PVD : FD->parameters()) {
- ImplicitParamDecl *IPD = ImplicitParamDecl::Create(
- getASTContext(), OFD, SourceLocation(), PVD->getIdentifier(),
- PVD->getType(), ImplicitParamKind::Other);
- OFD->setParam(i, IPD);
- ParmMap[PVD] = IPD;
- ++i;
- }
+ BuildSYCLKernelEntryPointOutline(SemaRef, FD, Body);
+ assert(OFD);
- OutlinedFunctionDeclBodyInstantiator OFDBodyInstantiator(SemaRef, ParmMap);
- Stmt *OFDBody = OFDBodyInstantiator.TransformStmt(Body).get();
- OFD->setBody(OFDBody);
- OFD->setNothrow();
- Stmt *NewBody = new (getASTContext()) SYCLKernelCallStmt(Body, OFD);
+ Stmt *NewBody =
+ new (getASTContext()) SYCLKernelCallStmt(Body, LaunchStmt, OFD);
return NewBody;
}
diff --git a/clang/lib/Serialization/ASTReaderStmt.cpp
b/clang/lib/Serialization/ASTReaderStmt.cpp
index 3f37dfbc3dea9..93276ce0b12ae 100644
--- a/clang/lib/Serialization/ASTReaderStmt.cpp
+++ b/clang/lib/Serialization/ASTReaderStmt.cpp
@@ -527,6 +527,7 @@ void ASTStmtReader::VisitCapturedStmt(CapturedStmt *S) {
void ASTStmtReader::VisitSYCLKernelCallStmt(SYCLKernelCallStmt *S) {
VisitStmt(S);
S->setOriginalStmt(cast<CompoundStmt>(Record.readSubStmt()));
+ S->setKernelLaunchStmt(cast<Stmt>(Record.readSubStmt()));
S->setOutlinedFunctionDecl(readDeclAs<OutlinedFunctionDecl>());
}
diff --git a/clang/lib/Serialization/ASTWriterStmt.cpp
b/clang/lib/Serialization/ASTWriterStmt.cpp
index be9bad9e96cc1..49b8ef178f93e 100644
--- a/clang/lib/Serialization/ASTWriterStmt.cpp
+++ b/clang/lib/Serialization/ASTWriterStmt.cpp
@@ -611,6 +611,7 @@ void ASTStmtWriter::VisitCapturedStmt(CapturedStmt *S) {
void ASTStmtWriter::VisitSYCLKernelCallStmt(SYCLKernelCallStmt *S) {
VisitStmt(S);
Record.AddStmt(S->getOriginalStmt());
+ Record.AddStmt(S->getKernelLaunchStmt());
Record.AddDeclRef(S->getOutlinedFunctionDecl());
Code = serialization::STMT_SYCLKERNELCALL;
diff --git a/clang/test/AST/ast-print-sycl-kernel-call.cpp
b/clang/test/AST/ast-print-sycl-kernel-call.cpp
new file mode 100644
index 0000000000000..2243ee024be1a
--- /dev/null
+++ b/clang/test/AST/ast-print-sycl-kernel-call.cpp
@@ -0,0 +1,22 @@
+// RUN: %clang_cc1 -fsycl-is-host -ast-print %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -ast-print %s -o - | FileCheck %s
+
+struct sycl_kernel_launcher {
+ template<typename KernelName, typename... Ts>
+ void sycl_enqueue_kernel_launch(const char *, Ts...) {}
+
+ template<typename KernelName, typename KernelType>
+ void kernel_entry_point(KernelType kernel) {
+ kernel();
+ }
+// CHECK: template <typename KernelName, typename KernelType> void
kernel_entry_point(KernelType kernel) {
+// CHECK-NEXT: kernel();
+// CHECK-NEXT: }
+// CHECK: template<> void kernel_entry_point<KN, (lambda at
{{.*}})>((lambda at {{.*}}) kernel) {
+// CHECK-NEXT: kernel();
+// CHECK-NEXT: }
+};
+
+void f(sycl_kernel_launcher skl) {
+ skl.kernel_entry_point<struct KN>([]{});
+}
diff --git a/clang/test/ASTSYCL/ast-dump-sycl-kernel-call-stmt.cpp
b/clang/test/ASTSYCL/ast-dump-sycl-kernel-call-stmt.cpp
index 27604e237adbb..10bd5e478f1ed 100644
--- a/clang/test/ASTSYCL/ast-dump-sycl-kernel-call-stmt.cpp
+++ b/clang/test/ASTSYCL/ast-dump-sycl-kernel-call-stmt.cpp
@@ -41,6 +41,13 @@ void skep1() {
// CHECK: |-FunctionDecl {{.*}} skep1 'void ()'
// CHECK-NEXT: | |-SYCLKernelCallStmt {{.*}}
// CHECK-NEXT: | | |-CompoundStmt {{.*}}
+// CHECK-NEXT: | | |-CompoundStmt {{.*}}
+// CHECK-NEXT: | | | |-DeclStmt {{.*}}
+// CHECK-NEXT: | | | | `-VarDecl {{.*}} kernel_name 'const char *' extern
+// CHECK-NEXT: | | | `-BinaryOperator {{.*}} 'const char *' lvalue '='
+// CHECK-NEXT: | | | |-DeclRefExpr {{.*}} 'const char *' lvalue Var {{.*}}
'kernel_name' 'const char *'
+// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'const char *'
<ArrayToPointerDecay>
+// CHECK-NEXT: | | | `-StringLiteral {{.*}} 'const char[14]' lvalue
"_ZTS2KNILi1EE"
// CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
// CHECK-NEXT: | | `-CompoundStmt {{.*}}
// CHECK-NEXT: | `-SYCLKernelEntryPointAttr {{.*}} KN<1>
@@ -77,6 +84,13 @@ void skep2<KN<2>>(K<2>);
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void () const' lvalue
CXXMethod {{.*}} 'operator()' 'void () const'
// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'const K<2>' lvalue <NoOp>
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'K<2>' lvalue ParmVar {{.*}}
'k' 'K<2>'
+// CHECK-NEXT: | | |-CompoundStmt {{.*}}
+// CHECK-NEXT: | | | |-DeclStmt {{.*}}
+// CHECK-NEXT: | | | | `-VarDecl {{.*}} kernel_name 'const char *' extern
+// CHECK-NEXT: | | | `-BinaryOperator {{.*}} 'const char *' lvalue '='
+// CHECK-NEXT: | | | |-DeclRefExpr {{.*}} 'const char *' lvalue Var {{.*}}
'kernel_name' 'const char *'
+// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'const char *'
<ArrayToPointerDecay>
+// CHECK-NEXT: | | | `-StringLiteral {{.*}} 'const char[14]' lvalue
"_ZTS2KNILi2EE"
// CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit used k 'K<2>'
// CHECK-NEXT: | | `-CompoundStmt {{.*}}
@@ -123,6 +137,13 @@ void skep3<KN<3>>(K<3> k) {
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void () const' lvalue CXXMethod
{{.*}} 'operator()' 'void () const'
// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'const K<3>' lvalue <NoOp>
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'K<3>' lvalue ParmVar {{.*}} 'k'
'K<3>'
+// CHECK-NEXT: | | |-CompoundStmt {{.*}}
+// CHECK-NEXT: | | | |-DeclStmt {{.*}}
+// CHECK-NEXT: | | | | `-VarDecl {{.*}} kernel_name 'const char *' extern
+// CHECK-NEXT: | | | `-BinaryOperator {{.*}} 'const char *' lvalue '='
+// CHECK-NEXT: | | | |-DeclRefExpr {{.*}} 'const char *' lvalue Var {{.*}}
'kernel_name' 'const char *'
+// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'const char *'
<ArrayToPointerDecay>
+// CHECK-NEXT: | | | `-StringLiteral {{.*}} 'const char[14]' lvalue
"_ZTS2KNILi3EE"
// CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit used k 'K<3>'
// CHECK-NEXT: | | `-CompoundStmt {{.*}}
@@ -152,6 +173,13 @@ void skep4(K<4> k, int p1, int p2) {
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'p1'
'int'
// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'p2'
'int'
+// CHECK-NEXT: | | |-CompoundStmt {{.*}}
+// CHECK-NEXT: | | | |-DeclStmt {{.*}}
+// CHECK-NEXT: | | | | `-VarDecl {{.*}} kernel_name 'const char *' extern
+// CHECK-NEXT: | | | `-BinaryOperator {{.*}} 'const char *' lvalue '='
+// CHECK-NEXT: | | | |-DeclRefExpr {{.*}} 'const char *' lvalue Var {{.*}}
'kernel_name' 'const char *'
+// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'const char *'
<ArrayToPointerDecay>
+// CHECK-NEXT: | | | `-StringLiteral {{.*}} 'const char[14]' lvalue
"_ZTS2KNILi4EE"
// CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit used k 'K<4>'
// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit used p1 'int'
@@ -182,7 +210,14 @@ void skep5(int unused1, K<5> k, int unused2, int p, int
unused3) {
// CHECK-NEXT: | |-ParmVarDecl {{.*}} unused3 'int'
// CHECK-NEXT: | |-SYCLKernelCallStmt {{.*}}
// CHECK-NEXT: | | |-CompoundStmt {{.*}}
-// CHECK: | | `-OutlinedFunctionDecl {{.*}}
+// CHECK: | | |-CompoundStmt {{.*}}
+// CHECK-NEXT: | | | |-DeclStmt {{.*}}
+// CHECK-NEXT: | | | | `-VarDecl {{.*}} kernel_name 'const char *' extern
+// CHECK-NEXT: | | | `-BinaryOperator {{.*}} 'const char *' lvalue '='
+// CHECK-NEXT: | | | |-DeclRefExpr {{.*}} 'const char *' lvalue Var {{.*}}
'kernel_name' 'const char *'
+// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'const char *'
<ArrayToPointerDecay>
+// CHECK-NEXT: | | | `-StringLiteral {{.*}} 'const char[14]' lvalue
"_ZTS2KNILi5EE"
+// CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit unused1 'int'
// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit used k 'K<5>'
// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit unused2 'int'
@@ -227,6 +262,13 @@ void skep6(const S6 &k) {
// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'void (*)() const'
<FunctionToPointerDecay>
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void () const' lvalue CXXMethod
{{.*}} 'operator()' 'void () const'
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'const S6' lvalue ParmVar {{.*}}
'k' 'const S6 &'
+// CHECK-NEXT: | | |-CompoundStmt {{.*}}
+// CHECK-NEXT: | | | |-DeclStmt {{.*}}
+// CHECK-NEXT: | | | | `-VarDecl {{.*}} kernel_name 'const char *' extern
+// CHECK-NEXT: | | | `-BinaryOperator {{.*}} 'const char *' lvalue '='
+// CHECK-NEXT: | | | |-DeclRefExpr {{.*}} 'const char *' lvalue Var {{.*}}
'kernel_name' 'const char *'
+// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'const char *'
<ArrayToPointerDecay>
+// CHECK-NEXT: | | | `-StringLiteral {{.*}} 'const char[14]' lvalue
"_ZTS2KNILi6EE"
// CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit used k 'const S6 &'
// CHECK-NEXT: | | `-CompoundStmt {{.*}}
@@ -260,6 +302,13 @@ void skep7(S7 k) {
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void () const' lvalue CXXMethod
{{.*}} 'operator()' 'void () const'
// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'const S7' lvalue <NoOp>
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'S7' lvalue ParmVar {{.*}} 'k'
'S7'
+// CHECK-NEXT: | | |-CompoundStmt {{.*}}
+// CHECK-NEXT: | | | |-DeclStmt {{.*}}
+// CHECK-NEXT: | | | | `-VarDecl {{.*}} kernel_name 'const char *' extern
+// CHECK-NEXT: | | | `-BinaryOperator {{.*}} 'const char *' lvalue '='
+// CHECK-NEXT: | | | |-DeclRefExpr {{.*}} 'const char *' lvalue Var {{.*}}
'kernel_name' 'const char *'
+// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'const char *'
<ArrayToPointerDecay>
+// CHECK-NEXT: | | | `-StringLiteral {{.*}} 'const char[14]' lvalue
"_ZTS2KNILi7EE"
// CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit used k 'S7'
// CHECK-NEXT: | | `-CompoundStmt {{.*}}
diff --git a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
index b5687523aee36..0884a291a54aa 100644
--- a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
+++ b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
@@ -55,8 +55,8 @@ int main() {
// Verify that SYCL kernel caller functions are not emitted during host
// compilation.
//
-// CHECK-HOST-NOT: _ZTS26single_purpose_kernel_name
-// CHECK-HOST-NOT: _ZTSZ4mainE18lambda_kernel_name
+// CHECK-HOST-NOT: define {{.*}} @_ZTS26single_purpose_kernel_name
+// CHECK-HOST-NOT: define {{.*}} @_ZTSZ4mainE18lambda_kernel_name
// Verify that sycl_kernel_entry_point attributed functions are not emitted
// during device compilation.
@@ -64,13 +64,13 @@ int main() {
// CHECK-DEVICE-NOT: single_purpose_kernel_task
// CHECK-DEVICE-NOT: kernel_single_task
-// Verify that no code is generated for the bodies of sycl_kernel_entry_point
-// attributed functions during host compilation. ODR-use of these functions may
-// require them to be emitted, but they have no effect if called.
+// Verify that kernel launch code is generated for sycl_kernel_entry_point
+// attributed functions during host compilation.
//
// CHECK-HOST-LINUX: define dso_local void
@_Z26single_purpose_kernel_task21single_purpose_kernel() #{{[0-9]+}} {
// CHECK-HOST-LINUX-NEXT: entry:
// CHECK-HOST-LINUX-NEXT: %kernelFunc = alloca
%struct.single_purpose_kernel, align 1
+// CHECK-HOST-LINUX-NEXT: store ptr @.str, ptr @kernel_name, align 8
// CHECK-HOST-LINUX-NEXT: ret void
// CHECK-HOST-LINUX-NEXT: }
//
@@ -79,6 +79,7 @@ int main() {
// CHECK-HOST-LINUX-NEXT: %kernelFunc = alloca %class.anon, align 4
// CHECK-HOST-LINUX-NEXT: %coerce.dive = getelementptr inbounds nuw
%class.anon, ptr %kernelFunc, i32 0, i32 0
// CHECK-HOST-LINUX-NEXT: store i32 %kernelFunc.coerce, ptr %coerce.dive,
align 4
+// CHECK-HOST-LINUX-NEXT: store ptr @.str.1, ptr @kernel_name, align 8
// CHECK-HOST-LINUX-NEXT: ret void
// CHECK-HOST-LINUX-NEXT: }
//
@@ -87,6 +88,7 @@ int main() {
// CHECK-HOST-WINDOWS-NEXT: %kernelFunc = alloca
%struct.single_purpose_kernel, align 1
// CHECK-HOST-WINDOWS-NEXT: %coerce.dive = getelementptr inbounds nuw
%struct.single_purpose_kernel, ptr %kernelFunc, i32 0, i32 0
// CHECK-HOST-WINDOWS-NEXT: store i8 %kernelFunc.coerce, ptr %coerce.dive,
align 1
+// CHECK-HOST-WINDOWS-NEXT: store ptr
@"??_C@_0CB@KFIJOMLB@_ZTS26single_purpose_kernel_name@", ptr
@"?kernel_name@?0??single_purpose_kernel_task@@YAXUsingle_purpose_kernel@@@Z@3PEBDEB",
align 8
// CHECK-HOST-WINDOWS-NEXT: ret void
// CHECK-HOST-WINDOWS-NEXT: }
//
@@ -95,6 +97,7 @@ int main() {
// CHECK-HOST-WINDOWS-NEXT: %kernelFunc = alloca %class.anon, align 4
// CHECK-HOST-WINDOWS-NEXT: %coerce.dive = getelementptr inbounds nuw
%class.anon, ptr %kernelFunc, i32 0, i32 0
// CHECK-HOST-WINDOWS-NEXT: store i32 %kernelFunc.coerce, ptr %coerce.dive,
align 4
+// CHECK-HOST-WINDOWS-NEXT: store ptr
@"??_C@_0BC@NHCDOLAA@_ZTSZ4mainEUlT_E_?$AA@", ptr
@"?kernel_name@?0???$kernel_single_task@V<lambda_1>@?0??main@@9@V1?0??2@9@@@YAXV<lambda_1>@?0??main@@9@@Z@3PEBDEB",
align 8
// CHECK-HOST-WINDOWS-NEXT: ret void
// CHECK-HOST-WINDOWS-NEXT: }
diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-appertainment.cpp
b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-appertainment.cpp
index 4774c8ef545f8..362dc06f1669e 100644
--- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-appertainment.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-appertainment.cpp
@@ -1,5 +1,8 @@
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only
-fsycl-is-host -fcxx-exceptions -verify %s
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only
-fsycl-is-device -verify %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++20 -fsyntax-only
-fsycl-is-host -fcxx-exceptions -verify %s
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++20 -fsyntax-only
-fsycl-is-device -verify %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++23 -fsyntax-only
-fsycl-is-host -fcxx-exceptions -verify %s
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++23 -fsyntax-only
-fsycl-is-device -verify %s
// These tests validate appertainment for the sycl_kernel_entry_point
attribute.
@@ -131,6 +134,16 @@ struct S15 {
static T ok15();
};
+struct S16 {
+ // Non-static member function declaration.
+ [[clang::sycl_kernel_entry_point(KN<16>)]]
+ void ok16();
+};
+
+#if __cplusplus >= 202302L
+auto ok17 = [] [[clang::sycl_kernel_entry_point(KN<17>)]] -> void {};
+#endif
+
////////////////////////////////////////////////////////////////////////////////
// Invalid declarations.
@@ -163,13 +176,6 @@ struct B2 {
static int bad2;
};
-struct B3 {
- // Non-static member function declaration.
- // expected-error@+1 {{the 'clang::sycl_kernel_entry_point' attribute cannot
be applied to a non-static member function}}
- [[clang::sycl_kernel_entry_point(BADKN<3>)]]
- void bad3();
-};
-
// expected-error@+1 {{'clang::sycl_kernel_entry_point' attribute only applies
to functions}}
namespace [[clang::sycl_kernel_entry_point(BADKN<4>)]] bad4 {}
@@ -244,13 +250,13 @@ void bad19() {
#endif
struct B20 {
- // expected-error@+1 {{the 'clang::sycl_kernel_entry_point' attribute cannot
be applied to a non-static member function}}
+ // expected-error@+1 {{the 'clang::sycl_kernel_entry_point' attribute cannot
be applied to a constructor}}
[[clang::sycl_kernel_entry_point(BADKN<20>)]]
B20();
};
struct B21 {
- // expected-error@+1 {{the 'clang::sycl_kernel_entry_point' attribute cannot
be applied to a non-static member function}}
+ // expected-error@+1 {{the 'clang::sycl_kernel_entry_point' attribute cannot
be applied to a destructor}}
[[clang::sycl_kernel_entry_point(BADKN<21>)]]
~B21();
};
@@ -337,11 +343,6 @@ struct B34 {
[[noreturn]] friend void bad34() {}
};
-#if __cplusplus >= 202302L
-// expected-error@+1 {{the 'clang::sycl_kernel_entry_point' attribute cannot
be applied to a non-static member function}}
-auto bad35 = [] [[clang::sycl_kernel_entry_point(BADKN<35>)]] -> void {};
-#endif
-
#if __cplusplus >= 202302L
// expected-error@+1 {{the 'clang::sycl_kernel_entry_point' attribute only
applies to functions with a non-deduced 'void' return type}}
auto bad36 = [] [[clang::sycl_kernel_entry_point(BADKN<36>)]] static {};
diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp
b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp
index 8f81fa218c171..fd1f00ae05d7a 100644
--- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp
@@ -1,4 +1,6 @@
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only
-fsycl-is-host -verify %s
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only
-fsycl-is-device -verify %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++20 -fsyntax-only
-fsycl-is-host -verify %s
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++20 -fsyntax-only
-fsycl-is-device -verify %s
// These tests validate parsing of the sycl_kernel_entry_point argument list
diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name.cpp
b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name.cpp
index c7b83932fefe6..5a3b43be66daf 100644
--- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name.cpp
@@ -1,4 +1,6 @@
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only
-fsycl-is-host -verify %s
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only
-fsycl-is-device -verify %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++20 -fsyntax-only
-fsycl-is-host -verify %s
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++20 -fsyntax-only
-fsycl-is-device -verify %s
// These tests validate that the kernel name type argument provided to the
diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-sfinae.cpp
b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-sfinae.cpp
index 4c61570419629..3689adaab9b5b 100644
--- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-sfinae.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-sfinae.cpp
@@ -1,4 +1,6 @@
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only
-fsycl-is-host -verify %s
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only
-fsycl-is-device -verify %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++20 -fsyntax-only
-fsycl-is-host -verify %s
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++20 -fsyntax-only
-fsycl-is-device -verify %s
// These tests are intended to validate that a sycl_kernel_entry_point
attribute
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
