[clang] [SYCL] SYCL host kernel launch support for the sycl_kernel_entry_point attribute. (PR #152403)
Tom Honermann via cfe-commits
cfe-commits at lists.llvm.org
Wed Oct 1 10:53:08 PDT 2025
https://github.com/tahonermann updated https://github.com/llvm/llvm-project/pull/152403
>From d0b53b9550d4332049b3b28ca649e2cf47a0cfbc Mon Sep 17 00:00:00 2001
From: Tom Honermann <tom.honermann at intel.com>
Date: Wed, 6 Aug 2025 19:26:54 -0700
Subject: [PATCH 1/6] [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 d9dc8290b0e49..7318e7640478f 100644
--- a/clang/include/clang/AST/ASTNodeTraverser.h
+++ b/clang/include/clang/AST/ASTNodeTraverser.h
@@ -849,8 +849,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 02581c8e73299..c8cc94fa1e86f 100644
--- a/clang/include/clang/AST/RecursiveASTVisitor.h
+++ b/clang/include/clang/AST/RecursiveASTVisitor.h
@@ -3036,6 +3036,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 2504841f6df33..4445a9094a9b7 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -520,13 +520,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,
@@ -538,7 +538,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.
@@ -553,39 +553,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.
@@ -600,7 +604,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.
@@ -609,17 +613,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++
@@ -633,31 +643,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.
@@ -671,24 +683,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 4bee813edf645..013f867cadd82 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12989,9 +12989,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"
diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp
index bbb957067c4c8..4d941550bc55f 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/clang/lib/AST/ASTContext.cpp
@@ -15118,6 +15118,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 afccba8778fd2..5272ecba7fc93 100644
--- a/clang/lib/AST/StmtPrinter.cpp
+++ b/clang/lib/AST/StmtPrinter.cpp
@@ -586,7 +586,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 031ef73214e76..9e26b145e0589 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 fc65199a0f154..79bb7231bd59f 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 160d7353cacd9..67663ddec4cd2 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -16352,19 +16352,19 @@ Decl *Sema::ActOnFinishFunctionBody(Decl *dcl, Stmt *Body, bool IsInstantiation,
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 94413b5b92d22..15e3a39c86427 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"
@@ -1251,6 +1252,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:
@@ -1431,7 +1441,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 2f97f6290f0e8..2c436e6388a36 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -315,43 +315,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();
}
@@ -389,6 +392,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
@@ -432,6 +496,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,
@@ -440,6 +534,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() &&
@@ -451,29 +550,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 8e8e03c2451a0..29461ef29a4af 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 cd1d4d801951d..729778aaee191 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 at _0CB@KFIJOMLB at _ZTS26single_purpose_kernel_name@", ptr @"?kernel_name@?0??single_purpose_kernel_task@@YAXUsingle_purpose_kernel@@@Z at 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 at _0BC@NHCDOLAA at _ZTSZ4mainEUlT_E_?$AA@", ptr @"?kernel_name@?0???$kernel_single_task at V<lambda_1>@?0??main@@9 at V1?0??2 at 9@@@YAXV<lambda_1>@?0??main@@9@@Z at 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 9aba284145fcb..3f07feb87c9a1 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 at +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 at +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 at +1 {{the 'clang::sycl_kernel_entry_point' attribute cannot be applied to a non-static member function}}
+ // expected-error at +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 at +1 {{the 'clang::sycl_kernel_entry_point' attribute cannot be applied to a non-static member function}}
+ // expected-error at +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 at +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 at +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
>From 98680049220f8f434dd8a5f561817ed8f126248a Mon Sep 17 00:00:00 2001
From: Tom Honermann <tom.honermann at intel.com>
Date: Sat, 16 Aug 2025 10:54:45 -0700
Subject: [PATCH 2/6] Add diagnostics for use of 'this' in a potentially
evaluated expression.
---
.../clang/Basic/DiagnosticSemaKinds.td | 3 +
clang/lib/Sema/SemaSYCL.cpp | 19 +-
.../sycl-kernel-entry-point-attr-this.cpp | 183 ++++++++++++++++++
3 files changed, 202 insertions(+), 3 deletions(-)
create mode 100644 clang/test/SemaSYCL/sycl-kernel-entry-point-attr-this.cpp
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 013f867cadd82..baf633b2f92cb 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12997,6 +12997,9 @@ def err_sycl_entry_point_invalid : Error<
def err_sycl_entry_point_invalid_redeclaration : Error<
"the %0 kernel name argument does not match prior"
" declaration%diff{: $ vs $|}1,2">;
+def err_sycl_entry_point_invalid_this : Error<
+ "'this' cannot be%select{| implicitly}0 used in a potentially evaluated"
+ " expression in the body of a function declared with the %1 attribute">;
def err_sycl_kernel_name_conflict : Error<
"the %0 kernel name argument conflicts with a previous declaration">;
def warn_sycl_kernel_name_not_a_class_type : Warning<
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 2c436e6388a36..90defb2826ea8 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -463,9 +463,10 @@ class OutlinedFunctionDeclBodyInstantiator
public:
using ParmDeclMap = llvm::DenseMap<ParmVarDecl *, VarDecl *>;
- OutlinedFunctionDeclBodyInstantiator(Sema &S, ParmDeclMap &M)
+ OutlinedFunctionDeclBodyInstantiator(Sema &S, ParmDeclMap &M,
+ FunctionDecl *FD)
: TreeTransform<OutlinedFunctionDeclBodyInstantiator>(S), SemaRef(S),
- MapRef(M) {}
+ MapRef(M), FD(FD) {}
// A new set of AST nodes is always required.
bool AlwaysRebuild() { return true; }
@@ -491,9 +492,20 @@ class OutlinedFunctionDeclBodyInstantiator
return DRE;
}
+ // Diagnose CXXThisExpr in a potentially evaluated expression.
+ ExprResult TransformCXXThisExpr(CXXThisExpr *CTE) {
+ if (SemaRef.currentEvaluationContext().isPotentiallyEvaluated()) {
+ SemaRef.Diag(CTE->getExprLoc(), diag::err_sycl_entry_point_invalid_this)
+ << (CTE->isImplicitCXXThis() ? /* implicit */ 1 : /* empty */ 0)
+ << FD->getAttr<SYCLKernelEntryPointAttr>();
+ }
+ return CTE;
+ }
+
private:
Sema &SemaRef;
ParmDeclMap &MapRef;
+ FunctionDecl *FD;
};
OutlinedFunctionDecl *BuildSYCLKernelEntryPointOutline(Sema &SemaRef,
@@ -518,7 +530,8 @@ OutlinedFunctionDecl *BuildSYCLKernelEntryPointOutline(Sema &SemaRef,
// evaluated contexts in the function body. This is not necessarily the
// right place to add such a diagnostic.
- OutlinedFunctionDeclBodyInstantiator OFDBodyInstantiator(SemaRef, ParmMap);
+ OutlinedFunctionDeclBodyInstantiator OFDBodyInstantiator(SemaRef, ParmMap,
+ FD);
Stmt *OFDBody = OFDBodyInstantiator.TransformStmt(Body).get();
OFD->setBody(OFDBody);
OFD->setNothrow();
diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-this.cpp b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-this.cpp
new file mode 100644
index 0000000000000..fc0640e1900cb
--- /dev/null
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-this.cpp
@@ -0,0 +1,183 @@
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -fsyntax-only -std=c++17 -fsycl-is-host -verify %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -fsyntax-only -std=c++17 -fsycl-is-device -verify %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -fsyntax-only -std=c++20 -fsycl-is-host -verify -DCXX20 %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -fsyntax-only -std=c++20 -fsycl-is-device -verify -DCXX20 %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -fsyntax-only -std=c++23 -fsycl-is-host -verify -DCXX23 %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -fsyntax-only -std=c++23 -fsycl-is-device -verify -DCXX23 %s
+
+// These tests validate diagnostics for invalid use of 'this' in the body of
+// a function declared with the sycl_kernel_entry_point attribute.
+
+
+template<typename T> struct remove_reference_t {
+ using type = T;
+};
+template<typename T> struct remove_reference_t<T&> {
+ using type = T;
+};
+
+namespace std {
+struct type_info {
+ virtual ~type_info();
+};
+} // namespace std
+
+////////////////////////////////////////////////////////////////////////////////
+// Valid declarations.
+////////////////////////////////////////////////////////////////////////////////
+template<int, int=0> struct KN;
+
+struct S1 {
+ [[clang::sycl_kernel_entry_point(KN<1>)]] void ok1() {
+ (void)sizeof(this);
+ }
+};
+
+struct S2 {
+ [[clang::sycl_kernel_entry_point(KN<2>)]] void ok2() {
+ (void)noexcept(this);
+ }
+};
+
+struct S3 {
+ [[clang::sycl_kernel_entry_point(KN<3>)]] void ok3() {
+ decltype(this) x = nullptr;
+ }
+};
+
+struct S4 {
+ static void smf();
+ [[clang::sycl_kernel_entry_point(KN<4>)]] void ok4() {
+ remove_reference_t<decltype(*this)>::type::smf();
+ }
+};
+
+struct S5 {
+ int dm;
+ void mf();
+ [[clang::sycl_kernel_entry_point(KN<5>)]] void ok5() {
+ (void)typeid(*this); // S5 is not abstract, so 'this' is not evaluated.
+ (void)typeid(dm); // 'int' is not an abstract class type; implicit 'this' is not evaluated.
+ (void)typeid(mf()); // 'void' is not an abstract class type; implicit 'this' is not evaluated.
+ }
+};
+
+template<typename KN, bool B>
+struct S6 {
+ void mf() noexcept(B);
+ [[clang::sycl_kernel_entry_point(KN)]] void ok6() noexcept(noexcept(mf())) {}
+};
+template void S6<KN<6,0>, false>::ok6();
+template void S6<KN<6,1>, true>::ok6();
+
+template<typename KN, bool B>
+struct S7 {
+ void mf() noexcept(B);
+ [[clang::sycl_kernel_entry_point(KN)]] void ok7() noexcept(noexcept(this->mf())) {}
+};
+template void S7<KN<7,0>, false>::ok7();
+template void S7<KN<7,1>, true>::ok7();
+
+#if defined(CXX20)
+template<typename KN, typename T>
+struct S8 {
+ void mf(T);
+ [[clang::sycl_kernel_entry_point(KN)]] void ok8() requires(requires { mf(1); }) {}
+};
+template void S8<KN<8>, int>::ok8();
+
+template<typename KN, typename T>
+struct S9 {
+ void mf(T);
+ [[clang::sycl_kernel_entry_point(KN)]] void ok9() requires(requires { this->mf(1); }) {}
+};
+template void S9<KN<9>, int>::ok9();
+#endif
+
+#if defined(CXX23)
+struct S10 {
+ [[clang::sycl_kernel_entry_point(KN<10>)]] void ok10(this S10 self) {
+ (void)self;
+ }
+};
+#endif
+
+
+////////////////////////////////////////////////////////////////////////////////
+// Invalid declarations.
+////////////////////////////////////////////////////////////////////////////////
+
+template<int, int=0> struct BADKN;
+
+// expected-error at +3 {{'this' cannot be used in a potentially evaluated expression in the body of a function declared with the 'clang::sycl_kernel_entry_point' attribute}}
+struct B1 {
+ [[clang::sycl_kernel_entry_point(BADKN<1>)]] void bad1() {
+ (void)this;
+ }
+};
+
+// expected-error at +4 {{'this' cannot be implicitly used in a potentially evaluated expression in the body of a function declared with the 'clang::sycl_kernel_entry_point' attribute}}
+struct B2 {
+ int dm;
+ [[clang::sycl_kernel_entry_point(BADKN<2>)]] void bad2() {
+ (void)dm;
+ }
+};
+
+// expected-error at +4 {{'this' cannot be implicitly used in a potentially evaluated expression in the body of a function declared with the 'clang::sycl_kernel_entry_point' attribute}}
+struct B3 {
+ void mf();
+ [[clang::sycl_kernel_entry_point(BADKN<3>)]] void bad3() {
+ (void)mf();
+ }
+};
+
+// expected-error at +4 {{'this' cannot be used in a potentially evaluated expression in the body of a function declared with the 'clang::sycl_kernel_entry_point' attribute}}
+struct B4 {
+ virtual void vmf() = 0;
+ [[clang::sycl_kernel_entry_point(BADKN<4>)]] void bad4() {
+ (void)typeid(*this); // B4 is abstract, so 'this' is evaluated.
+ }
+};
+
+// A diagnostic is not currently issued for uninstantiated definitions. In this
+// case, a declaration is instantiated, but a definition isn't. A diagnostic
+// will be issued if a definition is instantiated (as the next test exercises).
+struct B5 {
+ template<typename KN>
+ [[clang::sycl_kernel_entry_point(KN)]] void bad5() {
+ (void)this;
+ }
+};
+extern template void B5::bad5<BADKN<5>>();
+
+// expected-error at +4 {{'this' cannot be used in a potentially evaluated expression in the body of a function declared with the 'clang::sycl_kernel_entry_point' attribute}}
+struct B6 {
+ template<typename KN>
+ [[clang::sycl_kernel_entry_point(KN)]] void bad6() {
+ (void)this;
+ }
+};
+// expected-note at +1 {{in instantiation of function template specialization 'B6::bad6<BADKN<6>>' requested here}}
+template void B6::bad6<BADKN<6>>();
+
+// A diagnostic is not currently issued for uninstantiated definitions. In this
+// case, a declaration is instantiated, but a definition isn't. A diagnostic
+// will be issued if a definition is instantiated (as the next test exercises).
+template<typename KN>
+struct B7 {
+ [[clang::sycl_kernel_entry_point(KN)]] void bad7() {
+ (void)this;
+ }
+};
+extern template void B7<BADKN<7>>::bad7();
+
+// expected-error at +4 {{'this' cannot be used in a potentially evaluated expression in the body of a function declared with the 'clang::sycl_kernel_entry_point' attribute}}
+template<typename KN>
+struct B8 {
+ [[clang::sycl_kernel_entry_point(KN)]] void bad8() {
+ (void)this;
+ }
+};
+// expected-note at +1 {{in instantiation of member function 'B8<BADKN<8>>::bad8' requested here}}
+template void B8<BADKN<8>>::bad8();
>From 8e5e80842cee8557566481c7580741e201ae9cb6 Mon Sep 17 00:00:00 2001
From: Tom Honermann <tom.honermann at intel.com>
Date: Sat, 16 Aug 2025 11:01:00 -0700
Subject: [PATCH 3/6] Remove stale FIXME comment regarding diagnostics for
'this'.
---
clang/lib/Sema/SemaSYCL.cpp | 4 ----
1 file changed, 4 deletions(-)
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 90defb2826ea8..271dbe13e349e 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -526,10 +526,6 @@ OutlinedFunctionDecl *BuildSYCLKernelEntryPointOutline(Sema &SemaRef,
++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,
FD);
Stmt *OFDBody = OFDBodyInstantiator.TransformStmt(Body).get();
>From 70f34c38a7b8a232e186cb25b2191b1c54a30d37 Mon Sep 17 00:00:00 2001
From: Tom Honermann <tom.honermann at intel.com>
Date: Sun, 17 Aug 2025 18:32:54 -0700
Subject: [PATCH 4/6] Add tests for kernel name character encoding concerns.
---
clang/lib/AST/ASTContext.cpp | 4 --
clang/lib/Sema/SemaSYCL.cpp | 5 +-
.../ast-dump-sycl-kernel-call-stmt.cpp | 21 ++++++
.../CodeGenSYCL/kernel-caller-entry-point.cpp | 65 ++++++++++++++++++-
4 files changed, 86 insertions(+), 9 deletions(-)
diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp
index 4d941550bc55f..bbb957067c4c8 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/clang/lib/AST/ASTContext.cpp
@@ -15118,10 +15118,6 @@ 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/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 271dbe13e349e..56963fcb09853 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -397,10 +397,7 @@ CompoundStmt *BuildSYCLKernelLaunchStmt(Sema &SemaRef, FunctionDecl *FD,
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.
+ // Prepare a string literal that contains the kernel name.
QualType KernelNameCharTy = Ctx.CharTy.withConst();
llvm::APInt KernelNameSize(Ctx.getTypeSize(Ctx.getSizeType()),
KernelName.size() + 1);
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 29461ef29a4af..cdff3fc8821a1 100644
--- a/clang/test/ASTSYCL/ast-dump-sycl-kernel-call-stmt.cpp
+++ b/clang/test/ASTSYCL/ast-dump-sycl-kernel-call-stmt.cpp
@@ -319,6 +319,27 @@ void skep7(S7 k) {
// CHECK-NEXT: | | `-DeclRefExpr {{.*}} 'S7' lvalue ImplicitParam {{.*}} 'k' 'S7'
// CHECK-NEXT: | `-SYCLKernelEntryPointAttr {{.*}} KN<7>
+// Symbol names generated for the kernel entry point function should be
+// representable in the ordinary literal encoding even when the kernel name
+// type is named with esoteric characters.
+struct \u03b4\u03c4\u03c7; // Delta Tau Chi (δτχ)
+struct S8 {
+ void operator()() const;
+};
+[[clang::sycl_kernel_entry_point(\u03b4\u03c4\u03c7)]]
+void skep8(S8 k) {
+ k();
+}
+// CHECK: |-FunctionDecl {{.*}} skep8 'void (S8)'
+// CHECK-NEXT: | |-ParmVarDecl {{.*}} used k 'S8'
+// CHECK-NEXT: | |-SYCLKernelCallStmt {{.*}}
+// CHECK-NEXT: | | |-CompoundStmt {{.*}}
+// CHECK: | | |-CompoundStmt {{.*}}
+// CHECK: | | | `-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+// CHECK-NEXT: | | | `-StringLiteral {{.*}} 'const char[12]' lvalue "_ZTS6\316\264\317\204\317\207"
+// CHECK: | | `-OutlinedFunctionDecl {{.*}}
+// CHECK: | `-SYCLKernelEntryPointAttr {{.*}}
+
void the_end() {}
// CHECK: `-FunctionDecl {{.*}} the_end 'void ()'
diff --git a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
index 729778aaee191..e88e4b7cf3149 100644
--- a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
+++ b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
@@ -44,19 +44,24 @@ void kernel_single_task(KernelType kernelFunc) {
kernelFunc(42);
}
+// Exercise code gen with kernel name types named with esoteric characters.
+struct \u03b4\u03c4\u03c7; // Delta Tau Chi (δτχ)
+
int main() {
single_purpose_kernel obj;
single_purpose_kernel_task(obj);
int capture;
auto lambda = [=](auto) { (void) capture; };
kernel_single_task<decltype(lambda)>(lambda);
+ kernel_single_task<\u03b4\u03c4\u03c7>([](int){});
}
// Verify that SYCL kernel caller functions are not emitted during host
// compilation.
//
// CHECK-HOST-NOT: define {{.*}} @_ZTS26single_purpose_kernel_name
-// CHECK-HOST-NOT: define {{.*}} @_ZTSZ4mainE18lambda_kernel_name
+// CHECK-HOST-NOT: define {{.*}} @_ZTSZ4mainEUlT_E_
+// CHECK-HOST-NOT: define {{.*}} @"_ZTS6\CE\B4\CF\84\CF\87"
// Verify that sycl_kernel_entry_point attributed functions are not emitted
// during device compilation.
@@ -67,6 +72,10 @@ int main() {
// Verify that kernel launch code is generated for sycl_kernel_entry_point
// attributed functions during host compilation.
//
+// CHECK-HOST-LINUX: @.str = private unnamed_addr constant [33 x i8] c"_ZTS26single_purpose_kernel_name\00", align 1
+// CHECK-HOST-LINUX: @.str.1 = private unnamed_addr constant [18 x i8] c"_ZTSZ4mainEUlT_E_\00", align 1
+// CHECK-HOST-LINUX: @.str.2 = private unnamed_addr constant [12 x i8] c"_ZTS6\CE\B4\CF\84\CF\87\00", align 1
+//
// 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
@@ -83,6 +92,13 @@ int main() {
// CHECK-HOST-LINUX-NEXT: ret void
// CHECK-HOST-LINUX-NEXT: }
//
+// CHECK-HOST-LINUX: define internal void @"_Z18kernel_single_taskI6\CE\B4\CF\84\CF\87Z4mainEUliE_EvT0_"() #{{[0-9]+}} {
+// CHECK-HOST-LINUX-NEXT: entry:
+// CHECK-HOST-LINUX-NEXT: %kernelFunc = alloca %class.anon.0, align 1
+// CHECK-HOST-LINUX-NEXT: store ptr @.str.2, ptr @kernel_name, align 8
+// CHECK-HOST-LINUX-NEXT: ret void
+// CHECK-HOST-LINUX-NEXT: }
+//
// CHECK-HOST-WINDOWS: define dso_local void @"?single_purpose_kernel_task@@YAXUsingle_purpose_kernel@@@Z"(i8 %kernelFunc.coerce) #{{[0-9]+}} {
// CHECK-HOST-WINDOWS-NEXT: entry:
// CHECK-HOST-WINDOWS-NEXT: %kernelFunc = alloca %struct.single_purpose_kernel, align 1
@@ -100,6 +116,15 @@ int main() {
// CHECK-HOST-WINDOWS-NEXT: store ptr @"??_C at _0BC@NHCDOLAA at _ZTSZ4mainEUlT_E_?$AA@", ptr @"?kernel_name@?0???$kernel_single_task at V<lambda_1>@?0??main@@9 at V1?0??2 at 9@@@YAXV<lambda_1>@?0??main@@9@@Z at 3PEBDEB", align 8
// CHECK-HOST-WINDOWS-NEXT: ret void
// CHECK-HOST-WINDOWS-NEXT: }
+//
+// CHECK-HOST-WINDOWS: define internal void @"??$kernel_single_task at U\CE\B4\CF\84\CF\87@@V<lambda_2>@?0??main@@9@@@YAXV<lambda_2>@?0??main@@9@@Z"(i8 %kernelFunc.coerce) #{{[0-9]+}} {
+// CHECK-HOST-WINDOWS-NEXT: entry:
+// CHECK-HOST-WINDOWS-NEXT: %kernelFunc = alloca %class.anon.0, align 1
+// CHECK-HOST-WINDOWS-NEXT: %coerce.dive = getelementptr inbounds nuw %class.anon.0, 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 at _0M@BCGAEMBE at _ZTS6?N?$LE?O?$IE?O?$IH?$AA@", ptr @"?kernel_name@?0???$kernel_single_task at U\CE\B4\CF\84\CF\87@@V<lambda_2>@?0??main@@9@@@YAXV<lambda_2>@?0??main@@9@@Z at 3PEBDEB", align 8
+// CHECK-HOST-WINDOWS-NEXT: ret void
+// CHECK-HOST-WINDOWS-NEXT: }
// Verify that SYCL kernel caller functions are emitted for each device target.
//
@@ -182,6 +207,44 @@ int main() {
// CHECK-SPIR-NEXT: }
// CHECK-SPIR: define internal spir_func void @_ZZ4mainENKUlT_E_clIiEEDaS_
+// IR for the SYCL kernel caller function generated for kernel_single_task with
+// the Delta Tau Chi type as the SYCL kernel name type.
+//
+// CHECK-AMDGCN: Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
+// CHECK-AMDGCN-NEXT: define dso_local amdgpu_kernel void @"_ZTS6\CE\B4\CF\84\CF\87"
+// CHECK-AMDGCN-SAME: (ptr addrspace(4) noundef byref(%class.anon.0) align 1 %0) #[[AMDGCN_ATTR0]] {
+// CHECK-AMDGCN-NEXT: entry:
+// CHECK-AMDGCN-NEXT: %coerce = alloca %class.anon.0, align 1, addrspace(5)
+// CHECK-AMDGCN-NEXT: %kernelFunc = addrspacecast ptr addrspace(5) %coerce to ptr
+// CHECK-AMDGCN-NEXT: call void @llvm.memcpy.p0.p4.i64(ptr align 1 %kernelFunc, ptr addrspace(4) align 1 %0, i64 1, i1 false)
+// CHECK-AMDGCN-NEXT: call void @_ZZ4mainENKUliE_clEi
+// CHECK-AMDGCN-SAME: (ptr noundef nonnull align 1 dereferenceable(1) %kernelFunc, i32 noundef 42) #[[AMDGCN_ATTR1:[0-9]+]]
+// CHECK-AMDGCN-NEXT: ret void
+// CHECK-AMDGCN-NEXT: }
+// CHECK-AMDGCN: define internal void @_ZZ4mainENKUliE_clEi
+//
+// CHECK-NVPTX: Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
+// CHECK-NVPTX-NEXT: define dso_local ptx_kernel void @"_ZTS6\CE\B4\CF\84\CF\87"
+// CHECK-NVPTX-SAME: (ptr noundef byval(%class.anon.0) align 1 %kernelFunc) #[[NVPTX_ATTR0:[0-9]+]] {
+// CHECK-NVPTX-NEXT: entry:
+// CHECK-NVPTX-NEXT: call void @_ZZ4mainENKUliE_clEi
+// CHECK-NVPTX-SAME: (ptr noundef nonnull align 1 dereferenceable(1) %kernelFunc, i32 noundef 42) #[[NVPTX_ATTR1:[0-9]+]]
+// CHECK-NVPTX-NEXT: ret void
+// CHECK-NVPTX-NEXT: }
+// CHECK-NVPTX: define internal void @_ZZ4mainENKUliE_clEi
+//
+// CHECK-SPIR: Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
+// CHECK-SPIR-NEXT: define {{[a-z_ ]*}}spir_kernel void @"_ZTS6\CE\B4\CF\84\CF\87"
+// CHECK-SPIR-SAME: (ptr noundef byval(%class.anon.0) align 1 %kernelFunc) #[[SPIR_ATTR0:[0-9]+]] {
+// CHECK-SPIR-NEXT: entry:
+// CHECK-SPIR-NEXT: %kernelFunc.ascast = addrspacecast ptr %kernelFunc to ptr addrspace(4)
+// CHECK-SPIR-NEXT: call spir_func void @_ZZ4mainENKUliE_clEi
+// CHECK-SPIR-SAME: (ptr addrspace(4) noundef align 1 dereferenceable_or_null(1) %kernelFunc.ascast, i32 noundef 42) #[[SPIR_ATTR1:[0-9]+]]
+// CHECK-SPIR-NEXT: ret void
+// CHECK-SPIR-NEXT: }
+// CHECK-SPIR: define internal spir_func void @_ZZ4mainENKUliE_clEi
+
+
// CHECK-AMDGCN: #[[AMDGCN_ATTR0]] = { convergent mustprogress noinline norecurse nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
// CHECK-AMDGCN: #[[AMDGCN_ATTR1]] = { convergent nounwind }
//
>From dd8f6cf17754d8a2eb4911b7ca62637bcf7f9bc3 Mon Sep 17 00:00:00 2001
From: Mariya Podchishchaeva <mariya.podchishchaeva at intel.com>
Date: Mon, 29 Sep 2025 19:16:00 +0200
Subject: [PATCH 5/6] [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 <tom at honermann.net>
* 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 <tom at honermann.net>
* Remove array decay
* Add windows run line to the sema test
---------
Co-authored-by: Tom Honermann <tom at honermann.net>
---
clang/include/clang/AST/RecursiveASTVisitor.h | 7 +
clang/include/clang/AST/StmtSYCL.h | 53 +++++
clang/include/clang/Basic/AttrDocs.td | 14 +-
.../clang/Basic/DiagnosticSemaKinds.td | 9 +
clang/include/clang/Basic/StmtNodes.td | 1 +
clang/include/clang/Sema/ScopeInfo.h | 4 +
clang/include/clang/Sema/SemaSYCL.h | 6 +-
.../include/clang/Serialization/ASTBitCodes.h | 3 +
clang/lib/AST/ComputeDependence.cpp | 1 +
clang/lib/AST/StmtPrinter.cpp | 5 +
clang/lib/AST/StmtProfile.cpp | 5 +
clang/lib/CodeGen/CGStmt.cpp | 1 +
clang/lib/Sema/SemaDecl.cpp | 33 ++-
clang/lib/Sema/SemaExceptionSpec.cpp | 3 +
clang/lib/Sema/SemaSYCL.cpp | 173 ++++++++++-----
clang/lib/Sema/TreeTransform.h | 19 ++
clang/lib/Serialization/ASTReaderStmt.cpp | 12 +
clang/lib/Serialization/ASTWriterStmt.cpp | 10 +
clang/lib/StaticAnalyzer/Core/ExprEngine.cpp | 1 +
clang/test/AST/ast-print-sycl-kernel-call.cpp | 2 +-
.../ast-dump-sycl-kernel-call-stmt.cpp | 209 ++++++++++++++----
.../ast-dump-sycl-kernel-entry-point.cpp | 3 +
.../CodeGenSYCL/kernel-caller-entry-point.cpp | 91 +++++++-
.../unique_stable_name_windows_diff.cpp | 2 +
.../test/SemaSYCL/sycl-host-kernel-launch.cpp | 199 +++++++++++++++++
...-kernel-entry-point-attr-appertainment.cpp | 4 +
.../sycl-kernel-entry-point-attr-grammar.cpp | 4 +
...el-entry-point-attr-kernel-name-module.cpp | 6 +-
...ernel-entry-point-attr-kernel-name-pch.cpp | 9 +-
...cl-kernel-entry-point-attr-kernel-name.cpp | 6 +
.../sycl-kernel-entry-point-attr-sfinae.cpp | 5 +
.../sycl-kernel-entry-point-attr-this.cpp | 5 +
clang/tools/libclang/CXCursor.cpp | 1 +
33 files changed, 780 insertions(+), 126 deletions(-)
create mode 100644 clang/test/SemaSYCL/sycl-host-kernel-launch.cpp
diff --git a/clang/include/clang/AST/RecursiveASTVisitor.h b/clang/include/clang/AST/RecursiveASTVisitor.h
index c8cc94fa1e86f..2e44c1371cdc0 100644
--- a/clang/include/clang/AST/RecursiveASTVisitor.h
+++ b/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, {})
diff --git a/clang/include/clang/AST/StmtSYCL.h b/clang/include/clang/AST/StmtSYCL.h
index 70d8137992110..c8da4987321a9 100644
--- a/clang/include/clang/AST/StmtSYCL.h
+++ b/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
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index 4445a9094a9b7..7ebadfb25d9ca 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/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
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index baf633b2f92cb..1cd9803ad8741 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/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 "
diff --git a/clang/include/clang/Basic/StmtNodes.td b/clang/include/clang/Basic/StmtNodes.td
index c9c173f5c7469..1696f55b813ad 100644
--- a/clang/include/clang/Basic/StmtNodes.td
+++ b/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).
diff --git a/clang/include/clang/Sema/ScopeInfo.h b/clang/include/clang/Sema/ScopeInfo.h
index 4f4d38c961140..f334f58ebd0a7 100644
--- a/clang/include/clang/Sema/ScopeInfo.h
+++ b/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.
///
diff --git a/clang/include/clang/Sema/SemaSYCL.h b/clang/include/clang/Sema/SemaSYCL.h
index 7ae556da2bec1..76046b765c0d6 100644
--- a/clang/include/clang/Sema/SemaSYCL.h
+++ b/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
diff --git a/clang/include/clang/Serialization/ASTBitCodes.h b/clang/include/clang/Serialization/ASTBitCodes.h
index 441047d64f48c..0f170a40f05ca 100644
--- a/clang/include/clang/Serialization/ASTBitCodes.h
+++ b/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,
diff --git a/clang/lib/AST/ComputeDependence.cpp b/clang/lib/AST/ComputeDependence.cpp
index e0cf0deb12bd2..0192763c65f09 100644
--- a/clang/lib/AST/ComputeDependence.cpp
+++ b/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"
diff --git a/clang/lib/AST/StmtPrinter.cpp b/clang/lib/AST/StmtPrinter.cpp
index 5272ecba7fc93..9cc251a020096 100644
--- a/clang/lib/AST/StmtPrinter.cpp
+++ b/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());
}
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index 2035fa7635f2a..f1fd1c869cc52 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/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()));
diff --git a/clang/lib/CodeGen/CGStmt.cpp b/clang/lib/CodeGen/CGStmt.cpp
index 9e26b145e0589..4b87f1622d1b5 100644
--- a/clang/lib/CodeGen/CGStmt.cpp
+++ b/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)
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 67663ddec4cd2..2108e271bd430 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/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();
diff --git a/clang/lib/Sema/SemaExceptionSpec.cpp b/clang/lib/Sema/SemaExceptionSpec.cpp
index 15e3a39c86427..94eb3d0df3a4c 100644
--- a/clang/lib/Sema/SemaExceptionSpec.cpp
+++ b/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:
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 56963fcb09853..dd332f06303ef 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -390,53 +390,23 @@ void SemaSYCL::CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD) {
}
}
-namespace {
+ExprResult SemaSYCL::BuildSYCLKernelLaunchIdExpr(FunctionDecl *FD,
+ QualType KNT) {
-CompoundStmt *BuildSYCLKernelLaunchStmt(Sema &SemaRef, FunctionDecl *FD,
- const std::string &KernelName) {
ASTContext &Ctx = SemaRef.getASTContext();
- SmallVector<Stmt *> Stmts;
+ // Some routines need a valid source location to work correctly.
+ SourceLocation BodyLoc =
+ FD->getEndLoc().isValid() ? FD->getEndLoc() : FD->getLocation();
- // Prepare a string literal that contains the kernel name.
- 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);
+ IdentifierInfo &LaunchFooName =
+ Ctx.Idents.get("sycl_kernel_launch", tok::TokenKind::identifier);
// Perform overload resolution for a call to an accessible (member) function
- // template named 'sycl_enqueue_kernel_launch' from within the definition of
- // FD where:
+ // template named 'sycl_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.
+ // - 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
@@ -444,10 +414,108 @@ CompoundStmt *BuildSYCLKernelLaunchStmt(Sema &SemaRef, FunctionDecl *FD,
// - 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());
+ LookupResult Result(SemaRef, &LaunchFooName, BodyLoc,
+ Sema::LookupOrdinaryName);
+ CXXScopeSpec SS;
+ SemaRef.LookupTemplateName(Result, SemaRef.getCurScope(), SS,
+ /*ObjectType=*/QualType(),
+ /*EnteringContext=*/false, BodyLoc);
+
+ if (Result.empty() || Result.isAmbiguous()) {
+ SemaRef.Diag(BodyLoc, SemaRef.getLangOpts().SYCLIsHost
+ ? diag::err_sycl_host_no_launch_function
+ : diag::warn_sycl_device_no_host_launch_function);
+ SemaRef.Diag(BodyLoc, diag::note_sycl_host_launch_function);
+
+ return ExprError();
+ }
+
+ TemplateArgumentListInfo TALI{BodyLoc, BodyLoc};
+ TemplateArgument KNTA = TemplateArgument(KNT);
+ TemplateArgumentLoc TAL =
+ SemaRef.getTrivialTemplateArgumentLoc(KNTA, QualType(), BodyLoc);
+ TALI.addArgument(TAL);
+ ExprResult IdExpr;
+ if (SemaRef.isPotentialImplicitMemberAccess(SS, Result,
+ /*IsAddressOfOperand=*/false))
+ // BuildPossibleImplicitMemberExpr creates UnresolvedMemberExpr. Using it
+ // allows to pass implicit/explicit this argument automatically.
+ IdExpr = SemaRef.BuildPossibleImplicitMemberExpr(SS, BodyLoc, Result, &TALI,
+ SemaRef.getCurScope());
+ else
+ IdExpr = SemaRef.BuildTemplateIdExpr(SS, BodyLoc, Result,
+ /*RequiresADL=*/true, &TALI);
+
+ // Can happen if SKEP attributed function is a static member, but the launcher
+ // is a regular member. Perhaps emit a note saying that we're in host code
+ // synthesis.
+ if (IdExpr.isInvalid())
+ return ExprError();
+
+ return IdExpr;
+}
- return LaunchStmt;
+StmtResult SemaSYCL::BuildUnresolvedSYCLKernelCallStmt(CompoundStmt *CS,
+ Expr *IdExpr) {
+ return UnresolvedSYCLKernelCallStmt::Create(SemaRef.getASTContext(), CS,
+ IdExpr);
+}
+
+namespace {
+
+void PrepareKernelArgumentsForKernelLaunch(SmallVectorImpl<Expr *> &Args,
+ const SYCLKernelInfo *SKI,
+ Sema &SemaRef,
+ SourceLocation Loc) {
+ assert(SKI && "Need a kernel!");
+ ASTContext &Ctx = SemaRef.getASTContext();
+
+ // Prepare a string literal that contains the kernel name.
+ const std::string KernelName = SKI->GetKernelName();
+ 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);
+ Expr *KernelNameExpr =
+ StringLiteral::Create(Ctx, KernelName, StringLiteralKind::Ordinary,
+ /*Pascal*/ false, KernelNameArrayTy, Loc);
+ Args.push_back(KernelNameExpr);
+
+ // Right now we simply forward the arguments of the skep-attributed function.
+ // With decomposition present there can be another logic.
+ // Make sure to use CurContext to avoid diagnostics that we're using a
+ // variable coming from another context. The function should be the same as in
+ // the kernel info though.
+ auto *FD = cast<FunctionDecl>(SemaRef.CurContext);
+ assert(declaresSameEntity(FD, SKI->getKernelEntryPointDecl()));
+ for (ParmVarDecl *PVD : FD->parameters()) {
+ QualType ParamType = PVD->getOriginalType().getNonReferenceType();
+ Expr *DRE = SemaRef.BuildDeclRefExpr(PVD, ParamType, VK_LValue, Loc);
+ assert(DRE);
+ Args.push_back(DRE);
+ }
+}
+
+StmtResult BuildSYCLKernelLaunchStmt(Sema &SemaRef,
+ const SYCLKernelInfo *SKI,
+ Expr *IdExpr, SourceLocation Loc) {
+ SmallVector<Stmt *> Stmts;
+ assert(SKI && "Need a Kernel!");
+
+ if (IdExpr) {
+ llvm::SmallVector<Expr *, 12> Args;
+ PrepareKernelArgumentsForKernelLaunch(Args, SKI, SemaRef, Loc);
+ ExprResult LaunchResult =
+ SemaRef.BuildCallExpr(SemaRef.getCurScope(), IdExpr, Loc, Args, Loc);
+ if (LaunchResult.isInvalid())
+ return StmtError();
+
+ Stmts.push_back(LaunchResult.get());
+ }
+
+ return CompoundStmt::Create(SemaRef.getASTContext(), Stmts,
+ FPOptionsOverride(), Loc, Loc);
}
// The body of a function declared with the [[sycl_kernel_entry_point]]
@@ -535,11 +603,11 @@ OutlinedFunctionDecl *BuildSYCLKernelEntryPointOutline(Sema &SemaRef,
} // unnamed namespace
StmtResult SemaSYCL::BuildSYCLKernelCallStmt(FunctionDecl *FD,
- CompoundStmt *Body) {
+ CompoundStmt *Body,
+ Expr *LaunchIdExpr) {
assert(!FD->isInvalidDecl());
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.
@@ -557,18 +625,19 @@ StmtResult SemaSYCL::BuildSYCLKernelCallStmt(FunctionDecl *FD,
assert(declaresSameEntity(SKI.getKernelEntryPointDecl(), FD) &&
"SYCL kernel name conflict");
- // Build the kernel launch statement.
- Stmt *LaunchStmt =
- BuildSYCLKernelLaunchStmt(SemaRef, FD, SKI.GetKernelName());
- assert(LaunchStmt);
-
// Build the outline of the synthesized device entry point function.
OutlinedFunctionDecl *OFD =
BuildSYCLKernelEntryPointOutline(SemaRef, FD, Body);
assert(OFD);
+ // Build host kernel launch stmt.
+ SourceLocation BodyLoc =
+ FD->getEndLoc().isValid() ? FD->getEndLoc() : FD->getLocation();
+ StmtResult LaunchRes =
+ BuildSYCLKernelLaunchStmt(SemaRef, &SKI, LaunchIdExpr, BodyLoc);
+
Stmt *NewBody =
- new (getASTContext()) SYCLKernelCallStmt(Body, LaunchStmt, OFD);
+ new (getASTContext()) SYCLKernelCallStmt(Body, LaunchRes.get(), OFD);
return NewBody;
}
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 1d14ead778446..9932d46824d2d 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -12933,6 +12933,25 @@ ExprResult TreeTransform<Derived>::TransformSYCLUniqueStableNameExpr(
E->getLocation(), E->getLParenLocation(), E->getRParenLocation(), NewT);
}
+template <typename Derived>
+StmtResult TreeTransform<Derived>::TransformUnresolvedSYCLKernelCallStmt(
+ UnresolvedSYCLKernelCallStmt *S) {
+ ExprResult IdExpr = getDerived().TransformExpr(S->getKernelLaunchIdExpr());
+
+ if (IdExpr.isInvalid())
+ return StmtError();
+
+ StmtResult Body = getDerived().TransformStmt(S->getOriginalStmt());
+ if (Body.isInvalid())
+ return StmtError();
+
+ StmtResult SR = SemaRef.SYCL().BuildSYCLKernelCallStmt(
+ cast<FunctionDecl>(SemaRef.CurContext), cast<CompoundStmt>(Body.get()),
+ IdExpr.get());
+
+ return SR;
+}
+
template<typename Derived>
ExprResult
TreeTransform<Derived>::TransformPredefinedExpr(PredefinedExpr *E) {
diff --git a/clang/lib/Serialization/ASTReaderStmt.cpp b/clang/lib/Serialization/ASTReaderStmt.cpp
index 93276ce0b12ae..7bf0353797064 100644
--- a/clang/lib/Serialization/ASTReaderStmt.cpp
+++ b/clang/lib/Serialization/ASTReaderStmt.cpp
@@ -593,6 +593,14 @@ void ASTStmtReader::VisitSYCLUniqueStableNameExpr(SYCLUniqueStableNameExpr *E) {
E->setTypeSourceInfo(Record.readTypeSourceInfo());
}
+void ASTStmtReader::VisitUnresolvedSYCLKernelCallStmt(
+ UnresolvedSYCLKernelCallStmt *S) {
+ VisitStmt(S);
+
+ S->setOriginalStmt(cast<CompoundStmt>(Record.readSubStmt()));
+ S->setKernelLaunchIdExpr(Record.readExpr());
+}
+
void ASTStmtReader::VisitPredefinedExpr(PredefinedExpr *E) {
VisitExpr(E);
bool HasFunctionName = Record.readInt();
@@ -3163,6 +3171,10 @@ Stmt *ASTReader::ReadStmtFromStream(ModuleFile &F) {
S = SYCLUniqueStableNameExpr::CreateEmpty(Context);
break;
+ case STMT_UNRESOLVED_SYCL_KERNEL_CALL:
+ S = UnresolvedSYCLKernelCallStmt::CreateEmpty(Context);
+ break;
+
case EXPR_OPENACC_ASTERISK_SIZE:
S = OpenACCAsteriskSizeExpr::CreateEmpty(Context);
break;
diff --git a/clang/lib/Serialization/ASTWriterStmt.cpp b/clang/lib/Serialization/ASTWriterStmt.cpp
index 49b8ef178f93e..31d0e9796acf3 100644
--- a/clang/lib/Serialization/ASTWriterStmt.cpp
+++ b/clang/lib/Serialization/ASTWriterStmt.cpp
@@ -670,6 +670,16 @@ void ASTStmtWriter::VisitSYCLUniqueStableNameExpr(SYCLUniqueStableNameExpr *E) {
Code = serialization::EXPR_SYCL_UNIQUE_STABLE_NAME;
}
+void ASTStmtWriter::VisitUnresolvedSYCLKernelCallStmt(
+ UnresolvedSYCLKernelCallStmt *S) {
+ VisitStmt(S);
+
+ Record.AddStmt(S->getOriginalStmt());
+ Record.AddStmt(S->getKernelLaunchIdExpr());
+
+ Code = serialization::STMT_UNRESOLVED_SYCL_KERNEL_CALL;
+}
+
void ASTStmtWriter::VisitPredefinedExpr(PredefinedExpr *E) {
VisitExpr(E);
diff --git a/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp b/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
index c853c00019c10..ecaee9ee0094c 100644
--- a/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
+++ b/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
@@ -1824,6 +1824,7 @@ void ExprEngine::Visit(const Stmt *S, ExplodedNode *Pred,
case Stmt::OMPTargetParallelGenericLoopDirectiveClass:
case Stmt::CapturedStmtClass:
case Stmt::SYCLKernelCallStmtClass:
+ case Stmt::UnresolvedSYCLKernelCallStmtClass:
case Stmt::OpenACCComputeConstructClass:
case Stmt::OpenACCLoopConstructClass:
case Stmt::OpenACCCombinedConstructClass:
diff --git a/clang/test/AST/ast-print-sycl-kernel-call.cpp b/clang/test/AST/ast-print-sycl-kernel-call.cpp
index 2243ee024be1a..64c6624b768c9 100644
--- a/clang/test/AST/ast-print-sycl-kernel-call.cpp
+++ b/clang/test/AST/ast-print-sycl-kernel-call.cpp
@@ -3,7 +3,7 @@
struct sycl_kernel_launcher {
template<typename KernelName, typename... Ts>
- void sycl_enqueue_kernel_launch(const char *, Ts...) {}
+ void sycl_kernel_launch(const char *, Ts...) {}
template<typename KernelName, typename KernelType>
void kernel_entry_point(KernelType kernel) {
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 cdff3fc8821a1..8e9a6fe84961a 100644
--- a/clang/test/ASTSYCL/ast-dump-sycl-kernel-call-stmt.cpp
+++ b/clang/test/ASTSYCL/ast-dump-sycl-kernel-call-stmt.cpp
@@ -34,6 +34,8 @@ template<int> struct K {
void operator()(Ts...) const {}
};
+template <typename KernelName, typename... Ts>
+void sycl_kernel_launch(const char *, Ts...) {}
[[clang::sycl_kernel_entry_point(KN<1>)]]
void skep1() {
@@ -42,12 +44,11 @@ void skep1() {
// 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: | | | `-CallExpr {{.*}}
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'void (*)(const char *)' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (const char *)' lvalue Function {{.*}} 'sycl_kernel_launch' {{.*}}
// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
-// CHECK-NEXT: | | | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi1EE"
+// CHECK-NEXT: | | | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi1EE"
// CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
// CHECK-NEXT: | | `-CompoundStmt {{.*}}
// CHECK-NEXT: | `-SYCLKernelEntryPointAttr {{.*}} KN<1>
@@ -64,9 +65,10 @@ void skep2<KN<2>>(K<2>);
// CHECK-NEXT: | |-TemplateTypeParmDecl {{.*}} KT
// CHECK-NEXT: | |-FunctionDecl {{.*}} skep2 'void (KT)'
// CHECK-NEXT: | | |-ParmVarDecl {{.*}} k 'KT'
-// CHECK-NEXT: | | |-CompoundStmt {{.*}}
-// CHECK-NEXT: | | | `-CallExpr {{.*}} '<dependent type>'
-// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'KT' lvalue ParmVar {{.*}} 'k' 'KT'
+// CHECK-NEXT: | | |-UnresolvedSYCLKernelCallStmt {{.*}}
+// CHECK-NEXT: | | | `-CompoundStmt {{.*}}
+// CHECK-NEXT: | | | `-CallExpr {{.*}} '<dependent type>'
+// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'KT' lvalue ParmVar {{.*}} 'k' 'KT'
// CHECK-NEXT: | | `-SYCLKernelEntryPointAttr {{.*}} KNT
// CHECK-NEXT: | `-FunctionDecl {{.*}} skep2 'void (K<2>)' explicit_instantiation_definition
@@ -85,12 +87,14 @@ void skep2<KN<2>>(K<2>);
// 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: | | | `-CallExpr {{.*}} 'void'
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (const char *, K<2>)' lvalue Function {{.*}} 'sycl_kernel_launch' {{.*}}
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+// CHECK-NEXT: | | | | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi2EE"
+// CHECK-NEXT: | | | `-CXXConstructExpr {{.*}} 'K<2>' 'void (const K<2> &) noexcept'
+// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} <NoOp>
+// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'K<2>' lvalue ParmVar {{.*}} 'k' 'K<2>'
// CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit used k 'K<2>'
// CHECK-NEXT: | | `-CompoundStmt {{.*}}
@@ -116,9 +120,10 @@ void skep3<KN<3>>(K<3> k) {
// CHECK-NEXT: | |-TemplateTypeParmDecl {{.*}} KT
// CHECK-NEXT: | |-FunctionDecl {{.*}} skep3 'void (KT)'
// CHECK-NEXT: | | |-ParmVarDecl {{.*}} k 'KT'
-// CHECK-NEXT: | | |-CompoundStmt {{.*}}
-// CHECK-NEXT: | | | `-CallExpr {{.*}} '<dependent type>'
-// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'KT' lvalue ParmVar {{.*}} 'k' 'KT'
+// CHECK-NEXT: | | |-UnresolvedSYCLKernelCallStmt {{.*}}
+// CHECK-NEXT: | | | `-CompoundStmt {{.*}}
+// CHECK-NEXT: | | | `-CallExpr {{.*}} '<dependent type>'
+// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'KT' lvalue ParmVar {{.*}} 'k' 'KT'
// CHECK-NEXT: | | `-SYCLKernelEntryPointAttr {{.*}} KNT
// CHECK-NEXT: | `-Function {{.*}} 'skep3' 'void (K<3>)'
@@ -138,12 +143,14 @@ void skep3<KN<3>>(K<3> k) {
// 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: | | | `-CallExpr {{.*}}
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'void (*)(const char *, K<3>)' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (const char *, K<3>)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (const char *, K<3>)' {{.*}}
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+// CHECK-NEXT: | | | | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi3EE"
+// CHECK-NEXT: | | | `-CXXConstructExpr {{.*}} 'K<3>' 'void (const K<3> &) noexcept'
+// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'const K<3>' lvalue <NoOp>
+// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'K<3>' lvalue ParmVar {{.*}} 'k' 'K<3>'
// CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit used k 'K<3>'
// CHECK-NEXT: | | `-CompoundStmt {{.*}}
@@ -174,12 +181,18 @@ void skep4(K<4> k, int p1, int p2) {
// 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: | | | `-CallExpr {{.*}} 'void'
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'void (*)(const char *, K<4>, int, int)' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (const char *, K<4>, int, int)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (const char *, K<4>, int, int)' {{.*}}
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+// CHECK-NEXT: | | | | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi4EE"
+// CHECK-NEXT: | | | |-CXXConstructExpr {{.*}} 'K<4>' 'void (const K<4> &) noexcept'
+// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'const K<4>' lvalue <NoOp>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'K<4>' lvalue ParmVar {{.*}} 'k' 'K<4>'
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'p1' 'int'
+// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'p2' 'int'
// CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit used k 'K<4>'
// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit used p1 'int'
@@ -211,12 +224,22 @@ void skep5(int unused1, K<5> k, int unused2, int p, int unused3) {
// CHECK-NEXT: | |-SYCLKernelCallStmt {{.*}}
// CHECK-NEXT: | | |-CompoundStmt {{.*}}
// 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: | | | `-CallExpr {{.*}} 'void'
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'void (*)(const char *, int, K<5>, int, int, int)' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (const char *, int, K<5>, int, int, int)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (const char *, int, K<5>, int, int, int)' {{.*}}
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+// CHECK-NEXT: | | | | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi5EE"
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'unused1' 'int'
+// CHECK-NEXT: | | | |-CXXConstructExpr {{.*}} 'K<5>' 'void (const K<5> &) noexcept'
+// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'const K<5>' lvalue <NoOp>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'K<5>' lvalue ParmVar {{.*}} 'k' 'K<5>'
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'unused2' 'int'
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'p' 'int'
+// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'unused3' 'int'
// CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit unused1 'int'
// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit used k 'K<5>'
@@ -263,12 +286,13 @@ void skep6(const S6 &k) {
// 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: | | | `-CallExpr {{.*}} 'void'
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'void (*)(const char *, S6)' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (const char *, S6)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (const char *, S6)' {{.*}}
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+// CHECK-NEXT: | | | | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi6EE"
+// CHECK-NEXT: | | | `-CXXConstructExpr {{.*}} 'S6' 'void (const S6 &) noexcept'
+// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'const S6' lvalue ParmVar {{.*}} 'k' 'const S6 &'
// CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit used k 'const S6 &'
// CHECK-NEXT: | | `-CompoundStmt {{.*}}
@@ -303,12 +327,14 @@ void skep7(S7 k) {
// 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: | | | `-CallExpr {{.*}} 'void'
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'void (*)(const char *, S7)' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (const char *, S7)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (const char *, S7)' {{.*}}
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+// CHECK-NEXT: | | | | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi7EE"
+// CHECK-NEXT: | | | `-CXXConstructExpr {{.*}} 'S7' 'void (const S7 &) noexcept'
+// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'const S7' lvalue <NoOp>
+// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'S7' lvalue ParmVar {{.*}} 'k' 'S7'
// CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit used k 'S7'
// CHECK-NEXT: | | `-CompoundStmt {{.*}}
@@ -335,11 +361,96 @@ void skep8(S8 k) {
// CHECK-NEXT: | |-SYCLKernelCallStmt {{.*}}
// CHECK-NEXT: | | |-CompoundStmt {{.*}}
// CHECK: | | |-CompoundStmt {{.*}}
-// CHECK: | | | `-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
-// CHECK-NEXT: | | | `-StringLiteral {{.*}} 'const char[12]' lvalue "_ZTS6\316\264\317\204\317\207"
+// CHECK-NEXT: | | | `-CallExpr {{.*}} 'void'
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'void (*)(const char *, S8)' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (const char *, S8)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (const char *, S8)' {{.*}}
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+// CHECK-NEXT: | | | | `-StringLiteral {{.*}} 'const char[12]' lvalue "_ZTS6\316\264\317\204\317\207"
+// CHECK-NEXT: | | | `-CXXConstructExpr {{.*}} 'S8' 'void (const S8 &) noexcept'
+// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'const S8' lvalue <NoOp>
+// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'S8' lvalue ParmVar {{.*}} 'k' 'S8'
// CHECK: | | `-OutlinedFunctionDecl {{.*}}
// CHECK: | `-SYCLKernelEntryPointAttr {{.*}}
+class Handler {
+template <typename KernelName, typename... Ts>
+void sycl_kernel_launch(const char *, Ts...) {}
+public:
+template<typename KNT, typename KT>
+[[clang::sycl_kernel_entry_point(KNT)]]
+void skep9(KT k, int a, int b) {
+ k(a, b);
+}
+};
+void foo() {
+ Handler H;
+ H.skep9<KN<9>>([=](int a, int b){return a+b;}, 1, 2);
+}
+
+// CHECK: | |-FunctionTemplateDecl {{.*}} skep9
+// CHECK-NEXT: | | |-TemplateTypeParmDecl {{.*}} referenced typename depth 0 index 0 KNT
+// CHECK-NEXT: | | |-TemplateTypeParmDecl {{.*}} referenced typename depth 0 index 1 KT
+// CHECK-NEXT: | | |-CXXMethodDecl {{.*}} skep9 'void (KT, int, int)' implicit-inline
+// CHECK-NEXT: | | | |-ParmVarDecl {{.*}} referenced k 'KT'
+// CHECK-NEXT: | | | |-ParmVarDecl {{.*}} referenced a 'int'
+// CHECK-NEXT: | | | |-ParmVarDecl {{.*}} referenced b 'int'
+// CHECK-NEXT: | | | |-UnresolvedSYCLKernelCallStmt {{.*}}
+// CHECK-NEXT: | | | | `-CompoundStmt {{.*}}
+// CHECK-NEXT: | | | | `-CallExpr {{.*}} '<dependent type>'
+// CHECK-NEXT: | | | | |-DeclRefExpr {{.*}} 'KT' lvalue ParmVar {{.*}} 'k' 'KT'
+// CHECK-NEXT: | | | | |-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'a' 'int'
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'b' 'int'
+// CHECK-NEXT: | | | `-SYCLKernelEntryPointAttr {{.*}} KNT
+// CHECK-NEXT: | | `-CXXMethodDecl {{.*}} used skep9 {{.*}} implicit_instantiation implicit-inline
+// CHECK-NEXT: | | |-TemplateArgument type 'KN<9>'
+// CHECK-NEXT: | | | `-RecordType {{.*}} 'KN<9>' canonical
+// CHECK-NEXT: | | | `-ClassTemplateSpecialization {{.*}}'KN'
+// CHECK-NEXT: | | |-TemplateArgument type {{.*}}
+// CHECK-NEXT: | | | `-RecordType {{.*}}
+// CHECK-NEXT: | | | `-CXXRecord {{.*}}
+// CHECK-NEXT: | | |-ParmVarDecl {{.*}} used k {{.*}}
+// CHECK-NEXT: | | |-ParmVarDecl {{.*}} used a 'int'
+// CHECK-NEXT: | | |-ParmVarDecl {{.*}} used b 'int'
+// CHECK-NEXT: | | |-SYCLKernelCallStmt {{.*}}
+// CHECK-NEXT: | | | |-CompoundStmt {{.*}}
+// CHECK-NEXT: | | | | `-CXXOperatorCallExpr {{.*}} 'int' '()'
+// CHECK-NEXT: | | | | |-ImplicitCastExpr {{.*}} 'int (*)(int, int) const' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | | `-DeclRefExpr {{.*}} 'int (int, int) const' lvalue CXXMethod {{.*}} 'operator()' 'int (int, int) const'
+// CHECK-NEXT: | | | | |-ImplicitCastExpr {{.*}} lvalue <NoOp>
+// CHECK-NEXT: | | | | | `-DeclRefExpr {{.*}} lvalue ParmVar {{.*}} 'k' {{.*}}
+// CHECK-NEXT: | | | | |-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: | | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'a' 'int'
+// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'b' 'int'
+// CHECK-NEXT: | | | |-CompoundStmt {{.*}}
+// CHECK-NEXT: | | | | `-CXXMemberCallExpr {{.*}} 'void'
+// CHECK-NEXT: | | | | |-MemberExpr {{.*}} '<bound member function type>' ->sycl_kernel_launch {{.*}}
+// CHECK-NEXT: | | | | | `-CXXThisExpr {{.*}} 'Handler *' implicit this
+// CHECK-NEXT: | | | | |-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+// CHECK-NEXT: | | | | | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi9EE"
+// CHECK-NEXT: | | | | |-CXXConstructExpr {{.*}}
+// CHECK-NEXT: | | | | | `-ImplicitCastExpr {{.*}} lvalue <NoOp>
+// CHECK-NEXT: | | | | | `-DeclRefExpr {{.*}} lvalue ParmVar {{.*}} 'k' {{.*}}
+// CHECK-NEXT: | | | | |-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: | | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'a' 'int'
+// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'b' 'int'
+// CHECK-NEXT: | | | `-OutlinedFunctionDecl {{.*}}
+// CHECK-NEXT: | | | |-ImplicitParamDecl {{.*}} implicit used k {{.*}}
+// CHECK-NEXT: | | | |-ImplicitParamDecl {{.*}} implicit used a 'int'
+// CHECK-NEXT: | | | |-ImplicitParamDecl {{.*}} implicit used b 'int'
+// CHECK-NEXT: | | | `-CompoundStmt {{.*}}
+// CHECK-NEXT: | | | `-CXXOperatorCallExpr {{.*}} 'int' '()'
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'int (*)(int, int) const' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int (int, int) const' lvalue CXXMethod {{.*}} 'operator()' 'int (int, int) const'
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} lvalue <NoOp>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} lvalue ImplicitParam {{.*}} 'k' {{.*}}
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int' lvalue ImplicitParam {{.*}} 'a' 'int'
+// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'int' lvalue ImplicitParam {{.*}} 'b' 'int'
+// CHECK-NEXT: | | `-SYCLKernelEntryPointAttr {{.*}} struct KN<9>
+
void the_end() {}
// CHECK: `-FunctionDecl {{.*}} the_end 'void ()'
diff --git a/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp b/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp
index 1a82bdc1f5698..bfa3f764ceb0f 100644
--- a/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp
+++ b/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp
@@ -28,6 +28,9 @@
// A unique kernel name type is required for each declared kernel entry point.
template<int, int=0> struct KN;
+template <typename KernelName, typename... Tys>
+void sycl_kernel_launch(const char *, Tys &&...Args) {}
+
[[clang::sycl_kernel_entry_point(KN<1>)]]
void skep1() {
}
diff --git a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
index e88e4b7cf3149..48c13240b1ffc 100644
--- a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
+++ b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
@@ -25,8 +25,13 @@
// Test the generation of SYCL kernel caller functions. These functions are
// generated from functions declared with the sycl_kernel_entry_point attribute
-// and emited during device compilation. They are not emitted during device
-// compilation.
+// and emited during device compilation.
+// Test the generation of SYCL kernel launch statements during host compilation.
+// These statements are calls to sycl_enqueus_kernel_launch functions or class
+// members in case skep-attributed functions are also members of the same class.
+
+template <typename KernelName, typename KernelObj>
+void sycl_kernel_launch(const char *, KernelObj) {}
struct single_purpose_kernel_name;
struct single_purpose_kernel {
@@ -47,6 +52,17 @@ void kernel_single_task(KernelType kernelFunc) {
// Exercise code gen with kernel name types named with esoteric characters.
struct \u03b4\u03c4\u03c7; // Delta Tau Chi (δτχ)
+class Handler {
+template <typename KernelName, typename... Ts>
+void sycl_kernel_launch(const char *, Ts...) {}
+public:
+template<typename KNT, typename KT>
+[[clang::sycl_kernel_entry_point(KNT)]]
+void skep(KT k, int a, int b) {
+ k(a, b);
+}
+};
+
int main() {
single_purpose_kernel obj;
single_purpose_kernel_task(obj);
@@ -54,6 +70,8 @@ int main() {
auto lambda = [=](auto) { (void) capture; };
kernel_single_task<decltype(lambda)>(lambda);
kernel_single_task<\u03b4\u03c4\u03c7>([](int){});
+ Handler H;
+ H.skep<class notaverygoodkernelname>([=](int a, int b){return a+b;}, 1, 2);
}
// Verify that SYCL kernel caller functions are not emitted during host
@@ -79,50 +97,105 @@ int main() {
// 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: %agg.tmp = alloca %struct.single_purpose_kernel, align 1
+// CHECK-HOST-LINUX-NEXT: call void @_Z18sycl_kernel_launchI26single_purpose_kernel_name21single_purpose_kernelEvPKcT0_(ptr noundef @.str)
// CHECK-HOST-LINUX-NEXT: ret void
// CHECK-HOST-LINUX-NEXT: }
//
// CHECK-HOST-LINUX: define internal void @_Z18kernel_single_taskIZ4mainEUlT_E_S1_EvT0_(i32 %kernelFunc.coerce) #{{[0-9]+}} {
// CHECK-HOST-LINUX-NEXT: entry:
// CHECK-HOST-LINUX-NEXT: %kernelFunc = alloca %class.anon, align 4
+// CHECK-HOST-LINUX-NEXT: %agg.tmp = 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: call void @llvm.memcpy.p0.p0.i64(ptr align 4 %agg.tmp, ptr align 4 %kernelFunc, i64 4, i1 false)
+// CHECK-HOST-LINUX-NEXT: %coerce.dive1 = getelementptr inbounds nuw %class.anon, ptr %agg.tmp, i32 0, i32 0
+// CHECK-HOST-LINUX-NEXT: %0 = load i32, ptr %coerce.dive1, align 4
+// CHECK-HOST-LINUX-NEXT: call void @_Z18sycl_kernel_launchIZ4mainEUlT_E_S1_EvPKcT0_(ptr noundef @.str.1, i32 %0)
// CHECK-HOST-LINUX-NEXT: ret void
// CHECK-HOST-LINUX-NEXT: }
//
// CHECK-HOST-LINUX: define internal void @"_Z18kernel_single_taskI6\CE\B4\CF\84\CF\87Z4mainEUliE_EvT0_"() #{{[0-9]+}} {
// CHECK-HOST-LINUX-NEXT: entry:
// CHECK-HOST-LINUX-NEXT: %kernelFunc = alloca %class.anon.0, align 1
-// CHECK-HOST-LINUX-NEXT: store ptr @.str.2, ptr @kernel_name, align 8
+// CHECK-HOST-LINUX-NEXT: %agg.tmp = alloca %class.anon.0, align 1
+// CHECK-HOST-LINUX-NEXT: call void @"_Z18sycl_kernel_launchI6\CE\B4\CF\84\CF\87Z4mainEUliE_EvPKcT0_"(ptr noundef @.str.2)
// CHECK-HOST-LINUX-NEXT: ret void
// CHECK-HOST-LINUX-NEXT: }
-//
+
+// CHECK-HOST-LINUX: define internal void @_ZN7Handler4skepIZ4mainE22notaverygoodkernelnameZ4mainEUliiE_EEvT0_ii(ptr noundef nonnull align 1 dereferenceable(1) %this, i32 noundef %a, i32 noundef %b) #0 align 2 {
+// CHECK-HOST-LINUX-NEXT: entry:
+// CHECK-HOST-LINUX-NEXT: %k = alloca %class.anon.1, align 1
+// CHECK-HOST-LINUX-NEXT: %this.addr = alloca ptr, align 8
+// CHECK-HOST-LINUX-NEXT: %a.addr = alloca i32, align 4
+// CHECK-HOST-LINUX-NEXT: %b.addr = alloca i32, align 4
+// CHECK-HOST-LINUX-NEXT: %agg.tmp = alloca %class.anon.1, align 1
+// CHECK-HOST-LINUX-NEXT: store ptr %this, ptr %this.addr, align 8
+// CHECK-HOST-LINUX-NEXT: store i32 %a, ptr %a.addr, align 4
+// CHECK-HOST-LINUX-NEXT: store i32 %b, ptr %b.addr, align 4
+// CHECK-HOST-LINUX-NEXT: %this1 = load ptr, ptr %this.addr, align 8
+// CHECK-HOST-LINUX-NEXT: %0 = load i32, ptr %a.addr, align 4
+// CHECK-HOST-LINUX-NEXT: %1 = load i32, ptr %b.addr, align 4
+// CHECK-HOST-LINUX-NEXT: call void @_ZN7Handler18sycl_kernel_launchIZ4mainE22notaverygoodkernelnameJZ4mainEUliiE_iiEEEvPKcDpT0_(ptr noundef nonnull align 1 dereferenceable(1) %this1, ptr noundef @.str.3, i32 noundef %0, i32 noundef %1)
+// CHECK-HOST-LINUX-NEXT: ret void
+// CHECK-HOST-LINUX-NEXT: }
+
// CHECK-HOST-WINDOWS: define dso_local void @"?single_purpose_kernel_task@@YAXUsingle_purpose_kernel@@@Z"(i8 %kernelFunc.coerce) #{{[0-9]+}} {
// CHECK-HOST-WINDOWS-NEXT: entry:
// CHECK-HOST-WINDOWS-NEXT: %kernelFunc = alloca %struct.single_purpose_kernel, align 1
+// CHECK-HOST-WINDOWS-NEXT: %agg.tmp = 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 at _0CB@KFIJOMLB at _ZTS26single_purpose_kernel_name@", ptr @"?kernel_name@?0??single_purpose_kernel_task@@YAXUsingle_purpose_kernel@@@Z at 3PEBDEB", align 8
+// CHECK-HOST-WINDOWS-NEXT: %coerce.dive1 = getelementptr inbounds nuw %struct.single_purpose_kernel, ptr %agg.tmp, i32 0, i32 0
+// CHECK-HOST-WINDOWS-NEXT: %0 = load i8, ptr %coerce.dive1, align 1
+// CHECK-HOST-WINDOWS-NEXT: call void @"??$sycl_kernel_launch at Usingle_purpose_kernel_name@@Usingle_purpose_kernel@@@@YAXPEBDUsingle_purpose_kernel@@@Z"(ptr noundef @"??_C at _0CB@KFIJOMLB at _ZTS26single_purpose_kernel_name@", i8 %0)
// CHECK-HOST-WINDOWS-NEXT: ret void
// CHECK-HOST-WINDOWS-NEXT: }
//
// CHECK-HOST-WINDOWS: define internal void @"??$kernel_single_task at V<lambda_1>@?0??main@@9 at V1?0??2 at 9@@@YAXV<lambda_1>@?0??main@@9@@Z"(i32 %kernelFunc.coerce) #{{[0-9]+}} {
// CHECK-HOST-WINDOWS-NEXT: entry:
// CHECK-HOST-WINDOWS-NEXT: %kernelFunc = alloca %class.anon, align 4
+// CHECK-HOST-WINDOWS-NEXT: %agg.tmp = 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 at _0BC@NHCDOLAA at _ZTSZ4mainEUlT_E_?$AA@", ptr @"?kernel_name@?0???$kernel_single_task at V<lambda_1>@?0??main@@9 at V1?0??2 at 9@@@YAXV<lambda_1>@?0??main@@9@@Z at 3PEBDEB", align 8
+// CHECK-HOST-WINDOWS-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 4 %agg.tmp, ptr align 4 %kernelFunc, i64 4, i1 false)
+// CHECK-HOST-WINDOWS-NEXT: %coerce.dive1 = getelementptr inbounds nuw %class.anon, ptr %agg.tmp, i32 0, i32 0
+// CHECK-HOST-WINDOWS-NEXT: %0 = load i32, ptr %coerce.dive1, align 4
+// CHECK-HOST-WINDOWS-NEXT: call void @"??$sycl_kernel_launch at V<lambda_1>@?0??main@@9 at V1?0??2 at 9@@@YAXPEBDV<lambda_1>@?0??main@@9@@Z"(ptr noundef @"??_C at _0BC@NHCDOLAA at _ZTSZ4mainEUlT_E_?$AA@", i32 %0)
+//
// CHECK-HOST-WINDOWS-NEXT: ret void
// CHECK-HOST-WINDOWS-NEXT: }
//
// CHECK-HOST-WINDOWS: define internal void @"??$kernel_single_task at U\CE\B4\CF\84\CF\87@@V<lambda_2>@?0??main@@9@@@YAXV<lambda_2>@?0??main@@9@@Z"(i8 %kernelFunc.coerce) #{{[0-9]+}} {
// CHECK-HOST-WINDOWS-NEXT: entry:
// CHECK-HOST-WINDOWS-NEXT: %kernelFunc = alloca %class.anon.0, align 1
+// CHECK-HOST-WINDOWS-NEXT: %agg.tmp = alloca %class.anon.0, align 1
// CHECK-HOST-WINDOWS-NEXT: %coerce.dive = getelementptr inbounds nuw %class.anon.0, 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 at _0M@BCGAEMBE at _ZTS6?N?$LE?O?$IE?O?$IH?$AA@", ptr @"?kernel_name@?0???$kernel_single_task at U\CE\B4\CF\84\CF\87@@V<lambda_2>@?0??main@@9@@@YAXV<lambda_2>@?0??main@@9@@Z at 3PEBDEB", align 8
+// CHECK-HOST-WINDOWS-NEXT: %coerce.dive1 = getelementptr inbounds nuw %class.anon.0, ptr %agg.tmp, i32 0, i32 0
+// CHECK-HOST-WINDOWS-NEXT: %0 = load i8, ptr %coerce.dive1, align 1
+// CHECK-HOST-WINDOWS-NEXT: call void @"??$sycl_kernel_launch at U\CE\B4\CF\84\CF\87@@V<lambda_2>@?0??main@@9@@@YAXPEBDV<lambda_2>@?0??main@@9@@Z"(ptr noundef @"??_C at _0M@BCGAEMBE at _ZTS6?N?$LE?O?$IE?O?$IH?$AA@", i8 %0)
+// CHECK-HOST-WINDOWS-NEXT: ret void
+// CHECK-HOST-WINDOWS-NEXT: }
+
+// CHECK-HOST-WINDOWS: define internal void @"??$skep at Vnotaverygoodkernelname@?1??main@@9 at V<lambda_3>@?0??2 at 9@@Handler@@QEAAXV<lambda_3>@?0??main@@9 at HH@Z"(ptr noundef nonnull align 1 dereferenceable(1) %this, i8 %k.coerce, i32 noundef %a, i32 noundef %b) #0 align 2 {
+// CHECK-HOST-WINDOWS-NEXT: entry:
+// CHECK-HOST-WINDOWS-NEXT: %k = alloca %class.anon.1, align 1
+// CHECK-HOST-WINDOWS-NEXT: %b.addr = alloca i32, align 4
+// CHECK-HOST-WINDOWS-NEXT: %a.addr = alloca i32, align 4
+// CHECK-HOST-WINDOWS-NEXT: %this.addr = alloca ptr, align 8
+// CHECK-HOST-WINDOWS-NEXT: %agg.tmp = alloca %class.anon.1, align 1
+// CHECK-HOST-WINDOWS-NEXT: %coerce.dive = getelementptr inbounds nuw %class.anon.1, ptr %k, i32 0, i32 0
+// CHECK-HOST-WINDOWS-NEXT: store i8 %k.coerce, ptr %coerce.dive, align 1
+// CHECK-HOST-WINDOWS-NEXT: store i32 %b, ptr %b.addr, align 4
+// CHECK-HOST-WINDOWS-NEXT: store i32 %a, ptr %a.addr, align 4
+// CHECK-HOST-WINDOWS-NEXT: store ptr %this, ptr %this.addr, align 8
+// CHECK-HOST-WINDOWS-NEXT: %this1 = load ptr, ptr %this.addr, align 8
+// CHECK-HOST-WINDOWS-NEXT: %0 = load i32, ptr %b.addr, align 4
+// CHECK-HOST-WINDOWS-NEXT: %1 = load i32, ptr %a.addr, align 4
+// CHECK-HOST-WINDOWS-NEXT: %coerce.dive2 = getelementptr inbounds nuw %class.anon.1, ptr %agg.tmp, i32 0, i32 0
+// CHECK-HOST-WINDOWS-NEXT: %2 = load i8, ptr %coerce.dive2, align 1
+// CHECK-HOST-WINDOWS-NEXT: call void @"??$sycl_kernel_launch at Vnotaverygoodkernelname@?1??main@@9 at V<lambda_3>@?0??2 at 9@HH at Handler@@AEAAXPEBDV<lambda_3>@?0??main@@9 at HH@Z"(ptr noundef nonnull align 1 dereferenceable(1) %this1, ptr noundef @"??_C at _0CE@NJIGCEIA at _ZTSZ4mainE22notaverygoodkerneln@", i8 %2, i32 noundef %1, i32 noundef %0)
// CHECK-HOST-WINDOWS-NEXT: ret void
// CHECK-HOST-WINDOWS-NEXT: }
diff --git a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp
index 14366a092a1fe..63db83c02bbef 100644
--- a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp
+++ b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp
@@ -1,6 +1,8 @@
// RUN: %clang_cc1 -triple spir64-unknown-unknown -aux-triple x86_64-pc-windows-msvc -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s '-D$ADDRSPACE=addrspace(1) '
// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -fsycl-is-host -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s '-D$ADDRSPACE='
+template <typename KernelName, typename KernelObj>
+void sycl_kernel_launch(const char *, KernelObj) {}
template<typename KN, typename Func>
[[clang::sycl_kernel_entry_point(KN)]] void kernel(Func F){
diff --git a/clang/test/SemaSYCL/sycl-host-kernel-launch.cpp b/clang/test/SemaSYCL/sycl-host-kernel-launch.cpp
new file mode 100644
index 0000000000000..eda0c4da489a9
--- /dev/null
+++ b/clang/test/SemaSYCL/sycl-host-kernel-launch.cpp
@@ -0,0 +1,199 @@
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only -fsycl-is-host -fcxx-exceptions -verify=host,expected %s
+// RUN: %clang_cc1 -triple x86_64-windows-msvc -std=c++17 -fsyntax-only -fsycl-is-host -fcxx-exceptions -verify=host,expected %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only -fsycl-is-device -fcxx-exceptions -verify=device,expected %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++20 -fsyntax-only -fsycl-is-host -fcxx-exceptions -verify=host,expected %s
+
+// A unique kernel name type is required for each declared kernel entry point.
+template<int, int = 0> struct KN;
+
+[[clang::sycl_kernel_entry_point(KN<1>)]]
+void nolauncher() {}
+// host-error at -1 {{unable to find suitable 'sycl_kernel_launch' function for host code synthesis}}
+// device-warning at -2 {{unable to find suitable 'sycl_kernel_launch' function for host code synthesis}}
+// expected-note at -3 {{define 'sycl_kernel_launch' function template to fix}}
+
+void sycl_kernel_launch(const char *, int arg);
+// expected-note at -1 {{declared as a non-template here}}
+
+[[clang::sycl_kernel_entry_point(KN<2>)]]
+void nontemplatel() {}
+// host-error at -1 {{unable to find suitable 'sycl_kernel_launch' function for host code synthesis}}
+// device-warning at -2 {{unable to find suitable 'sycl_kernel_launch' function for host code synthesis}}
+// expected-note at -3 {{define 'sycl_kernel_launch' function template to fix}}
+// expected-error at -4 {{'sycl_kernel_launch' following the 'template' keyword does not refer to a template}}
+
+template <typename KernName>
+void sycl_kernel_launch(const char *, int arg);
+// expected-note at -1 {{candidate function template not viable: requires 2 arguments, but 1 was provided}}
+// expected-note at -2 2{{candidate function template not viable: no known conversion from 'Kern' to 'int' for 2nd argument}}
+
+[[clang::sycl_kernel_entry_point(KN<3>)]]
+void notenoughargs() {}
+// expected-error at -1 {{no matching function for call to 'sycl_kernel_launch'}}
+// FIXME: Should this also say "no suitable function for host code synthesis"?
+
+
+template <typename KernName>
+void sycl_kernel_launch(const char *, bool arg = 1);
+// expected-note at -1 2{{candidate function template not viable: no known conversion from 'Kern' to 'bool' for 2nd argument}}
+
+[[clang::sycl_kernel_entry_point(KN<4>)]]
+void enoughargs() {}
+
+namespace boop {
+template <typename KernName, typename KernelObj>
+void sycl_kernel_launch(const char *, KernelObj);
+
+template <typename KernName, typename KernelObj>
+[[clang::sycl_kernel_entry_point(KernName)]]
+void iboop(KernelObj Kernel) {
+ Kernel();
+}
+}
+
+template <typename KernName, typename KernelObj>
+[[clang::sycl_kernel_entry_point(KernName)]]
+void idontboop(KernelObj Kernel) {
+ Kernel();
+}
+// expected-error at -3 {{no matching function for call to 'sycl_kernel_launch'}}
+
+struct Kern {
+ int a;
+ int *b;
+ Kern(int _a, int* _b) : a(_a), b(_b) {}
+ void operator()(){ *b = a;}
+};
+
+void foo() {
+ int *a;
+ Kern b(1, a);
+ idontboop<KN<6>>(b);
+ // expected-note at -1 {{in instantiation of function template specialization 'idontboop<KN<6>, Kern>' requested here}}
+ boop::iboop<KN<7>>(b);
+}
+
+class MaybeHandler {
+
+template <typename KernName>
+void sycl_kernel_launch(const char *);
+
+template <typename KernName, typename... Tys>
+void sycl_kernel_launch(const char *, Tys ...Args);
+
+public:
+
+template <typename KernName, typename KernelObj>
+[[clang::sycl_kernel_entry_point(KernName)]]
+void entry(KernelObj Kernel) {
+ Kernel();
+}
+};
+
+class MaybeHandler2 {
+
+template <typename KernName, typename... Tys>
+static void sycl_kernel_launch(const char *, Tys ...Args);
+
+public:
+
+template <typename KernName, typename KernelObj>
+[[clang::sycl_kernel_entry_point(KernName)]]
+void entry(KernelObj Kernel) {
+ Kernel();
+}
+};
+
+class MaybeHandler3 {
+
+template <typename KernName, typename... Tys>
+static void sycl_kernel_launch(const char *, Tys ...Args);
+
+public:
+
+template <typename KernName, typename KernelObj>
+[[clang::sycl_kernel_entry_point(KernName)]]
+static void entry(KernelObj Kernel) {
+ Kernel();
+}
+};
+
+class MaybeHandler4 {
+
+template <typename KernName, typename... Tys>
+void sycl_kernel_launch(const char *, Tys ...Args);
+
+public:
+
+template <typename KernName, typename KernelObj>
+[[clang::sycl_kernel_entry_point(KernName)]]
+static void entry(KernelObj Kernel) {
+ // expected-error at -1 {{call to non-static member function without an object argument}}
+ // FIXME: Should that be clearer?
+ Kernel();
+}
+};
+
+template<typename>
+struct base_handler {
+ template<typename KNT, typename... Ts>
+ void sycl_kernel_launch(const char*, Ts...) {}
+};
+struct derived_handler : base_handler<derived_handler> {
+ template<typename KNT, typename KT>
+ [[clang::sycl_kernel_entry_point(KNT)]]
+ void entry(KT k) { k(); }
+};
+
+template<int N>
+struct derived_handler_t : base_handler<derived_handler_t<N>> {
+ template<typename KNT, typename KT>
+// FIXME this fails because accessing members of dependent bases requires
+// explicit qualification.
+ [[clang::sycl_kernel_entry_point(KNT)]]
+ void entry(KT k) { k(); }
+ // expected-error at -1 {{no matching function for call to 'sycl_kernel_launch'}}
+};
+
+template<typename KNT>
+struct kernel_launcher {
+ template<typename... Ts>
+ void operator()(const char*, Ts...) const {}
+};
+
+namespace var {
+template<typename KNT>
+kernel_launcher<KNT> sycl_kernel_launch;
+
+struct handler {
+ template<typename KNT, typename KT>
+ [[clang::sycl_kernel_entry_point(KNT)]]
+ void entry(KT k) { k(); }
+};
+}
+
+
+void bar() {
+ int *a;
+ Kern b(1, a);
+ MaybeHandler H;
+ MaybeHandler2 H1;
+ MaybeHandler3 H2;
+ MaybeHandler4 H3;
+ H.entry<KN<8>>(b);
+ H1.entry<KN<9>>(b);
+ H2.entry<KN<10>>(b);
+ H3.entry<KN<11>>(b);
+
+ derived_handler H5;
+ H5.entry<KN<12>>(b);
+
+ derived_handler_t<13> H6;
+ H6.entry<KN<13>>(b); //expected-note {{in instantiation of function template specialization}}
+
+ var::handler h;
+ h.entry<KN<14>>(b);
+}
+
+
+
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 3f07feb87c9a1..c181f76321a26 100644
--- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-appertainment.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-appertainment.cpp
@@ -40,6 +40,10 @@ struct coroutine_traits {
// A unique kernel name type is required for each declared kernel entry point.
template<int, int = 0> struct KN;
+// A launcher function definition required for host code synthesis to silence
+// complains.
+template <typename KernelName, typename... Tys>
+void sycl_kernel_launch(const char *, Tys &&...Args) {}
////////////////////////////////////////////////////////////////////////////////
// Valid declarations.
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 fd1f00ae05d7a..1cdd48f1e5840 100644
--- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp
@@ -10,6 +10,10 @@
template<int> struct ST; // #ST-decl
template<int N> using TTA = ST<N>; // #TTA-decl
+// A launcher function definition required for host code synthesis to silence
+// complains.
+template <typename KernelName, typename... Tys>
+void sycl_kernel_launch(const char *, Tys &&...Args) {}
////////////////////////////////////////////////////////////////////////////////
// Valid declarations.
diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name-module.cpp b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name-module.cpp
index 8788e147a2ae4..44a3ce6f3640a 100644
--- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name-module.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name-module.cpp
@@ -17,6 +17,11 @@ module M2 { header "m2.h" }
#--- common.h
template<int> struct KN;
+// A launcher function definition required for host code synthesis to silence
+// complains.
+template <typename KernelName, typename... Tys>
+void sycl_kernel_launch(const char *, Tys &&...Args) {}
+
[[clang::sycl_kernel_entry_point(KN<1>)]]
void common_test1() {}
@@ -25,7 +30,6 @@ template<typename T>
void common_test2() {}
template void common_test2<KN<2>>();
-
#--- m1.h
#include "common.h"
diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name-pch.cpp b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name-pch.cpp
index 0575a7a5a67eb..0e6d1a6c57e39 100644
--- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name-pch.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name-pch.cpp
@@ -15,6 +15,11 @@
#--- pch.h
template<int> struct KN;
+// A launcher function definition required for host code synthesis to silence
+// complains.
+template <typename KernelName, typename... Tys>
+void sycl_kernel_launch(const char *, Tys &&...Args) {}
+
[[clang::sycl_kernel_entry_point(KN<1>)]]
void pch_test1() {} // << expected previous declaration note here.
@@ -26,11 +31,11 @@ template void pch_test2<KN<2>>();
#--- test.cpp
// expected-error at +3 {{the 'clang::sycl_kernel_entry_point' kernel name argument conflicts with a previous declaration}}
-// expected-note at pch.h:4 {{previous declaration is here}}
+// expected-note at pch.h:9 {{previous declaration is here}}
[[clang::sycl_kernel_entry_point(KN<1>)]]
void test1() {}
// expected-error at +3 {{the 'clang::sycl_kernel_entry_point' kernel name argument conflicts with a previous declaration}}
-// expected-note at pch.h:8 {{previous declaration is here}}
+// expected-note at pch.h:13 {{previous declaration is here}}
[[clang::sycl_kernel_entry_point(KN<2>)]]
void test2() {}
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 5a3b43be66daf..7b525abadd2c3 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
@@ -9,6 +9,12 @@
// specification.
struct S1;
+
+// A launcher function definition required for host code synthesis to silence
+// complains.
+template <typename KernelName, typename... Tys>
+void sycl_kernel_launch(const char *, Tys &&...Args) {}
+
// expected-warning at +3 {{redundant 'clang::sycl_kernel_entry_point' attribute}}
// expected-note at +1 {{previous attribute is here}}
[[clang::sycl_kernel_entry_point(S1),
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 3689adaab9b5b..9674dac456f9f 100644
--- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-sfinae.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-sfinae.cpp
@@ -10,6 +10,11 @@
// attribute during instantiation of a specialization unless that specialization
// is selected by overload resolution.
+// A launcher function definition required for host code synthesis to silence
+// complains.
+template <typename KernelName, typename... Tys>
+void sycl_kernel_launch(const char *, Tys &&...Args) {}
+
// FIXME: C++23 [temp.expl.spec]p12 states:
// FIXME: ... Similarly, attributes appearing in the declaration of a template
// FIXME: have no effect on an explicit specialization of that template.
diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-this.cpp b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-this.cpp
index fc0640e1900cb..7b8fc6c9a7630 100644
--- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-this.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-this.cpp
@@ -22,6 +22,11 @@ struct type_info {
};
} // namespace std
+// A launcher function definition required for host code synthesis to silence
+// complains.
+template <typename KernelName, typename... Tys>
+void sycl_kernel_launch(const char *, Tys &&...Args) {}
+
////////////////////////////////////////////////////////////////////////////////
// Valid declarations.
////////////////////////////////////////////////////////////////////////////////
diff --git a/clang/tools/libclang/CXCursor.cpp b/clang/tools/libclang/CXCursor.cpp
index 3c4062410eac1..ad44aea6e3f5e 100644
--- a/clang/tools/libclang/CXCursor.cpp
+++ b/clang/tools/libclang/CXCursor.cpp
@@ -376,6 +376,7 @@ CXCursor cxcursor::MakeCXCursor(const Stmt *S, const Decl *Parent,
break;
case Stmt::SYCLKernelCallStmtClass:
+ case Stmt::UnresolvedSYCLKernelCallStmtClass:
K = CXCursor_UnexposedStmt;
break;
>From 1271ba3c8ab57674742b018f9ba56593f167e755 Mon Sep 17 00:00:00 2001
From: Mariya Podchishchaeva <mariya.podchishchaeva at intel.com>
Date: Wed, 1 Oct 2025 19:52:56 +0200
Subject: [PATCH 6/6] [SYCL-Upstreaming] Fix a crash (#52)
In case a function with skep attribute is instantiated two times with
the same kernel name the attribute is invalid due to the conflicting name.
Make sure to exit from instantiation of UnresolvedSYCLKernelCallStmt in
this case.
---
clang/lib/Sema/TreeTransform.h | 5 +++++
...cl-kernel-entry-point-attr-kernel-name.cpp | 22 +++++++++++++++++++
2 files changed, 27 insertions(+)
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 9932d46824d2d..bf994b01fcb78 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -12936,6 +12936,11 @@ ExprResult TreeTransform<Derived>::TransformSYCLUniqueStableNameExpr(
template <typename Derived>
StmtResult TreeTransform<Derived>::TransformUnresolvedSYCLKernelCallStmt(
UnresolvedSYCLKernelCallStmt *S) {
+ auto *FD = cast<FunctionDecl>(SemaRef.CurContext);
+ const auto *SKEPAttr = FD->getAttr<SYCLKernelEntryPointAttr>();
+ if (!SKEPAttr || SKEPAttr->isInvalidAttr())
+ return StmtError();
+
ExprResult IdExpr = getDerived().TransformExpr(S->getKernelLaunchIdExpr());
if (IdExpr.isInvalid())
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 7b525abadd2c3..e2e6bf3314614 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
@@ -124,3 +124,25 @@ struct B19 {
};
// expected-note at +1 {{in instantiation of template class 'B19<int>' requested here}}
B19<int> b19;
+
+struct auto_name;
+
+// expected-error at +4 {{the 'clang::sycl_kernel_entry_point' kernel name argument conflicts with a previous declaration}}
+// expected-note at +3 {{previous declaration is here}}
+template <typename KernelName, typename KernelType>
+[[clang::sycl_kernel_entry_point(KernelName)]]
+void __kernel_single_task(const KernelType KernelFunc) {
+ KernelFunc();
+}
+
+template <typename KernelType, typename KernelName = auto_name>
+void pf(KernelType K) {
+ // expected-note at +1 {{requested here}}
+ __kernel_single_task<KernelName>(K);
+}
+
+void foo() {
+ pf([](){});
+ // expected-note at +1 {{requested here}}
+ pf([](){});
+}
More information about the cfe-commits
mailing list